/*
Copyright (c) 2010, Advanced Micro Devices
All rights reserved.

Redistribution and use in source and binary forms, with or without modification, are permitted 
provided that the following conditions are met:

Redistributions of source code must retain the above copyright notice, this list of conditions 
and the following disclaimer.

Redistributions in binary form must reproduce the above copyright notice, this list of 
conditions and the following disclaimer in the documentation and/or other materials provided 
with the distribution.

Neither the name of Advanced Micro Devices nor the names of its contributors may be used to endorse 
or promote products derived from this software without specific prior written permission.

THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR 
IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY 
AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR 
CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR 
CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR 
SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY 
THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR 
OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE 
POSSIBILITY OF SUCH DAMAGE.
*/

/*! \file
*   \brief PipelineKit backend code generator
*   \author Alan Heirich
*   
*   \version 0.1
*   \date May 2010
*/

#include "pipeline-kit.hpp"

#include <Set>

namespace amd{

	// Attribute class

	void
		Attribute::selfDeclare( ofstream * classFile )
	{
		*classFile << "double " << name() << ";" << std::endl;
	}

	Maybe<ASTExpr>::MaybeType 
		AttributeContainer::maybeExpr( std::string name )
	{
		if( attributes().find( name ) != attributes().end() )
		{
			Attribute * a = attributes()[ name ];
			return Maybe<ASTExpr>::MaybeType( a->value() );
		}
		if( pipeline()->attributes().find( name ) != pipeline()->attributes().end() )
		{
			Attribute * a = pipeline()->attributes()[ name ];
			return Maybe<ASTExpr>::MaybeType( a->value() );
		}
		if( pipeline()->equations().find( name ) != pipeline()->equations().end() ){
			Equation * equation = pipeline()->equations()[ name ];
			return Maybe<ASTExpr>::MaybeType( equation->value() );
		}
		return Maybe<ASTExpr>::MaybeType( boost::tuple<string>( "undefined" ) );
	}

	bool 
		AttributeContainer::isDefined( std::string name )
		{
			if( attributes().find( name ) != attributes().end() )
			{
				return true;
			}
			if( pipeline()->attributes().find( name ) != pipeline()->attributes().end() )
			{
				return true;
			}
			if( pipeline()->equations().find( name ) != pipeline()->equations().end() ){
				return true;
			}
			return false;
	}


	// Pipeline class

	Equation *
		Pipeline::declareEquation( 
		std::string name,
		std::string type,
		ASTExpr & value )
	{
		Equation * E = new Equation( name, type, value, this );
		equations()[ name ] = E;
		return E;
	}


	Code * 
		Pipeline::declareCode( 
		CODE_SECTION codeType, 
		std::string text,
		ASTAttributeList & attributes )
	{
		Code * C = new Code( codeType, text, attributes, this );
		codes().insert( pair<std::string, Code *>( C->name(), C ));
		return C;
	}


	Buffer * 
		Pipeline::declareBuffer( 
		std::string name, 
		ADDRESS_SPACE addressSpace, 
		std::string type, 
		ASTExpr & size, 
		ASTAttributeList & attributes )
	{
		Buffer * B = new Buffer( name, addressSpace, type, size, attributes, this );
		buffers()[ name ] = B;
		return B;
	}


	Image2D * 
		Pipeline::declareImage2D( 
		std::string name, 
		ADDRESS_SPACE addressSpace, 
		std::string type, 
		ASTExpr & sizeX, 
		ASTExpr & sizeY,
		ASTAttributeList & attributes )
	{
		Image2D * I = new Image2D( name, addressSpace, type, sizeX, sizeY, attributes, this ); 
		image2Ds()[ name ] = I;
		return I;
	}



	Image3D * 
		Pipeline::declareImage3D( 
		std::string name, 
		ADDRESS_SPACE addressSpace, 
		std::string type, 
		ASTExpr & sizeX, 
		ASTExpr & sizeY,
		ASTExpr & sizeZ,
		ASTAttributeList & attributes )
	{
		Image3D * I = new Image3D( name, addressSpace, type, sizeX, sizeY, sizeZ, attributes, this );
		image3Ds()[ name ] = I;
		return I;
	}


	Uniform * 
		Pipeline::declareUniform( 
		std::string name, 
		std::string type,
		ASTAttributeList & attributes )
	{
		Uniform * U = new Uniform( name, type, attributes, this );
		uniforms()[ name ] = U;
		return U;
	}


	Kernel *
		Pipeline::declareKernel( 
		std::string name, 
#if 1
		ASTTypeVarList & arguments,
#else
		ASTTypeVarList & inputs, 
		ASTTypeVarList & outputs, 
#endif
		ASTAttributeList & attributes )
	{
		Kernel * K = new Kernel( name, arguments, /*inputs, outputs,*/ attributes, this );
		kernels()[ name ] = K;
		return K;
	}


	Stage * 
		Pipeline::declareStage( 
		std::string name, 
		std::string kernel,
#if 1
		NameList & args,
#else
		NameList & inputs, 
		NameList & outputs, 
#endif
		ASTAttributeList & attributes )
	{
		Stage * S = new Stage( name, kernel, args, /* inputs, outputs,*/ attributes, this );
		stages()[ name ] = S;
		return S;
	}


	Graph * 
		Pipeline::declareGraph( 
		std::string name, 
		ASTGraphEdgeList & edges, 
		ASTAttributeList & attributes )
	{
		Graph * G = new Graph( name, edges, attributes, this );
		graphs()[ name ] = G;
		return G;
	}


	void
		Pipeline::importEquations( ASTVarDefMap environment )
	{
		ASTVarDefMap::iterator vdmIt;
		vdmIt = environment.begin();
		while( vdmIt != environment.end() ){
			std::string varName = vdmIt->first;
			ASTVarDef varDef = vdmIt->second;
			std::string varType = boost::get<0>( varDef );
			ASTExpr a = boost::get<1>( varDef );
			declareEquation( varName, varType, a );
			vdmIt++;
		}//while
	}


	// Class Code

	// Class Equation

	void
		Equation::selfDeclare( ofstream * classFile )
	{
		*classFile << type() << " " << name() << ";" << std::endl;
	}


	void
		Equation::selfDeclareAccessorMutator( ofstream * classFile )
	{
		*classFile << std::endl;
		*classFile << type() << std::endl;
		*classFile << "get" << name() << "( void ){" << std::endl;
		*classFile << "return " << name() << ";" << std::endl;
		*classFile << "}" << std::endl;

		*classFile << std::endl;
		*classFile << "void" << std::endl;
		*classFile << "set" << name() << "( " << type() << " value ){" << std::endl;
		*classFile << name() << " = value;" << std::endl;
		generateDependantEvaluations( classFile );
		*classFile << "}" << std::endl;
	}


	void
		Equation::generateDependantEvaluations( ofstream * classFile )
	{
		// be sure to invoke every equation that depends on this one
		EquationMap::iterator eqIt;
		eqIt = pipeline()->equations().begin();
		while( eqIt != pipeline()->equations().end() ){
			Equation * e = eqIt->second;
			if( e != this ){
				ASTExpr v = e->value();
				NameList precedents = expr2NameList( v );
				bool isDependant = false;
				NameList::iterator nIt;
				nIt = precedents.begin();
				while( nIt != precedents.end() ){
					std::string precedentName = *nIt;
					if( 0 == precedentName.compare( name() ) ){
						isDependant = true;
						break;
					}
					nIt++;
				}//while
				if( isDependant ){
					*classFile << "evaluate" << e->name() << "();" << std::endl;
				}
			}
			eqIt++;
		}//while
	}


	void
		Equation::generatePredecessorEvaluations( ofstream * classFile )
	{
		ASTExpr v = value();
		NameList precedents = expr2NameList( v );
		NameList::iterator namIt;
		namIt = precedents.begin();
		while( namIt != precedents.end() ){
			std::string precedentName = *namIt;
			if( pipeline()->equations().find( precedentName ) != pipeline()->equations().end() ){
				Equation * e = pipeline()->equations()[ precedentName ];
				if( !e->evaluated() ){
					e->generatePredecessorEvaluations( classFile );
					e->selfInvokeEvaluation( classFile );
				}
			}
			namIt++;
		}//while
	}

	void
		Equation::generateSelfEvaluationAssignment( ofstream * classFile )
	{
		ASTExpr v = value();
		if( stringify( v ).compare( "undefined" ) ){
			if( type().compare( "string" )){
				*classFile << name() << " = (" << type() << ")( " 
					<< stringify( v ) << " );" << std::endl;
			}
		}
	}

	void
		Equation::generateSelfEvaluation( ofstream * classFile )
	{
		UniformMap::iterator unIt =	pipeline()->uniforms().find( name() );
		if( unIt == pipeline()->uniforms().end() ){
			return;
		}
		Uniform * uniform = unIt->second;
		if( uniform->type().compare( "string" ) == 0 ){
			return;
		}
		*classFile << std::endl;
		*classFile << type() << " evaluate" << name() << "( void ){" << std::endl;

		generateSelfEvaluationAssignment( classFile );	
		generateDependantEvaluations( classFile );

		*classFile << "return " << name() << ";" << std::endl;
		*classFile << "}" << std::endl;
	}

	void
		Equation::selfInvokeEvaluation( ofstream * classFile )
	{
		*classFile << "evaluate" << name() <<"();" << std::endl;
		evaluatedIs( true );
	}


	bool
		Equation::allPrecedentsEvaluated()
	{
		ASTExpr v = value();
		NameList precedents = expr2NameList( v );
		if( precedents.size() == 0 ){
			return true;
		}
		NameList::iterator nIt;
		nIt = precedents.begin();
		while( nIt != precedents.end() ){
			std::string name = *nIt;
			nIt++;
			if( !name.compare( "undefined" )){
				undefinedIs( true );
				return false;
			}
			EquationMap::iterator eqIt;
			eqIt = pipeline()->equations().find( name );
			if( eqIt != pipeline()->equations().end() ){
				Equation * e = eqIt->second;
				if( e->undefined() ){
					undefinedIs( true );
					return false;
				}
				if( !e->evaluated() ){
					return false;
				}
				continue;
			}
			else{
				if( pipeline()->isDefined( name ) ){
					continue;
				}
				assert( "undefined reference" );
			}
		}//while
		return true;
	}

	void
		Equation::check()
	{
		pipeline()->checkName( name() );
		if( type().compare( "string" )){
			pipeline()->checkExpression( name(), value() );
		}
	}

	// Class Buffer

	Buffer::Buffer(
		std::string name,
		ADDRESS_SPACE addressSpace, 
		std::string type, 
		ASTExpr size, 
		ASTAttributeList attributes,
		Pipeline * pipeline )
	{
		set( name, addressSpace, type, size, attributes, pipeline );
	}


	void 
		Buffer::set(
		std::string name,
		ADDRESS_SPACE addressSpace, 
		std::string type, 
		ASTExpr size, 
		ASTAttributeList attributes,
		Pipeline * pipeline )
	{
		pipelineIs( pipeline );
		name_ = name;
		addressSpace_ = addressSpace;
		type_ = type;
		size_ = size;
		importAttributes( attributes );
	}

	void
		Buffer::selfDeclare( ofstream * classFile )
	{
		if( addressSpace() == AST_GLOBAL && pipeline()->targetIsOpenCL() ){
			if( isTrue( "glInterop" ) ){
				*classFile << "#ifdef GL_PIPELINE" << std::endl;
				*classFile << "cl::" << typeName() << "GL " << name() << ";" << std::endl;
				*classFile << "#endif" << std::endl;
			} 
			if( isTrue( "dxInterop" ) ){
				*classFile << "#ifdef DX_PIPELINE" << std::endl;
				*classFile << "cl::" << typeName() << "D3D10 " << name() << ";" << std::endl;
				*classFile << "#endif" << std::endl;
			}
			if( isTrue( "glInterop" ) || isTrue( "dxInterop" ) ){
				*classFile << "#if !defined( GL_PIPELINE ) && !defined( DX_PIPELINE )" << std::endl;
			}
			*classFile << "cl::" << typeName() << " " << name() << ";" << std::endl;
			if( isTrue( "glInterop" ) || isTrue( "dxInterop" ) ){
				*classFile << "#endif" << std::endl;
			}		
		}
		else if( pipeline()->targetIsDirectCompute() ){
			if( addressSpace() == AST_GLOBAL ){
				*classFile << "// " << name() << std::endl;
				*classFile << "ID3D11Buffer * " << name() << ";" << std::endl;
				*classFile << "ID3D11Buffer * " << name() << "Mappable;" << std::endl;
				*classFile << "ID3D11ShaderResourceView * " << SRVname() << ";" << std::endl;
				*classFile << "ID3D11UnorderedAccessView * " << UAVname() << ";" << std::endl;
			}
		}

	}


	void
		Buffer::selfDeclareDirectComputeBuffer( 
		ofstream * kernelFile, 
		Kernel * k, 
		int & idUAV,
		int & idSRV
		)
	{	
		if( addressSpace() == AST_GLOBAL ){
			if( isUsedAsOutput( k )){
				*kernelFile << "RWStructuredBuffer< " << type() << " > " 
					<< name() << " : register( u" << idUAV++ << " );" << std::endl;
			}
			else{
				*kernelFile << "StructuredBuffer< " << type() << " > " 
					<< name() << " : register( t" << idSRV++ << " );" << std::endl;
			}
		}
		else if( addressSpace() == AST_LOCAL ){
			// to do - does not work for computeShaders with Uniform-dependent array size
			// if sizeString() refers to any Uniforms or Equations then you need
			// to pass these in also
			*kernelFile << "groupshared "  << type() << " " << name() << "[ "
				<< sizeString() << " ];" << std::endl;
		}
	}

	void 
		Buffer::selfInitialize( ofstream * classFile )
	{
		if( addressSpace() == AST_GLOBAL ){
			if( pipeline()->targetIsOpenCL() ){
				std::string accessClass;
				if( isInput() && isOutput() ){
					accessClass = "CL_MEM_READ_WRITE";
				}
				else if( isInput() && !isOutput() ){
					accessClass = "CL_MEM_READ_ONLY";
				}
				else if( !isInput() && isOutput() ){
					accessClass = "CL_MEM_WRITE_ONLY";
				}
				else{
					accessClass = "CL_MEM_READ_WRITE";
				}
				*classFile << "{" << std::endl;
				// to do examine all buffers and identify as read only, write only, or read write
				if( isTrue( "glInterop" ) ){
					*classFile << "#if defined( GL_PIPELINE )" << std::endl;
					*classFile << name() << " = cl::" << typeName() << "GL( context, " << accessClass << ", (int)VBO( \""
						<< name() << "\" ));" << std::endl;
					*classFile << "#endif" << std::endl;
				}
				if( isTrue( "dxInterop" ) ){
					*classFile << "#if defined( DX_PIPELINE )" << std::endl;
					*classFile << name() << " = cl::" << typeName() << "D3D10( context, " << accessClass << ", (ID3D10Buffer *)VBO( \""
						<< name() << "\" ));" << std::endl;
					*classFile << "#endif" << std::endl;
				}
				bool useInterop = isTrue( "glInterop" ) || isTrue( "dxInterop" );
				if( useInterop ){
					*classFile << "#if !defined( GL_PIPELINE ) && !defined( DX_PIPELINE )" << std::endl;
				}
				*classFile << imageFormat() << std::endl;
				*classFile << name() << " = cl::" << typeName() << "( context, " << accessClass << ", ";
				if( isImage() ){
					*classFile << "cif, " << std::endl;
					*classFile << dimensions();
					*classFile << ", NULL, NULL ";
				}
				else{
					*classFile << dimensionsSizeT();
				}
				//to do instead of comparing to imageFormat make a bool function to indicate buffer versus image
				*classFile << " );" << std::endl;				
				if( useInterop ){
					*classFile << "#endif" << std::endl;
				}
				*classFile << "}" << std::endl;
			}
			else if( pipeline()->targetIsDirectCompute() ){
				*classFile << "{" << std::endl;
				*classFile << "// buffer " << name() << std::endl;
				*classFile << "D3D11_BUFFER_DESC bufferDesc;" << std::endl;
				*classFile << "memset( &bufferDesc, 0, sizeof( bufferDesc ) );" << std::endl;
				*classFile << "bufferDesc.ByteWidth = " << byteWidthString() << ";" << std::endl;
				*classFile << "bufferDesc.Usage = D3D11_USAGE_STAGING;" << std::endl;
				*classFile << "bufferDesc.CPUAccessFlags = D3D11_CPU_ACCESS_READ | D3D11_CPU_ACCESS_WRITE;" << std::endl;
				*classFile << "bufferDesc.StructureByteStride = " << baseCount() 
					<< " * sizeof( " << baseType() << ");" << std::endl;
				*classFile << "m_pd3dDevice->CreateBuffer( &bufferDesc, NULL, &" << name()
					<< "Mappable );" << std::endl;

				*classFile << std::endl;
				*classFile << "memset( &bufferDesc, 0, sizeof( bufferDesc ) );" << std::endl;
				*classFile << "bufferDesc.ByteWidth = " << byteWidthString() << ";" << std::endl;
				*classFile << "bufferDesc.Usage = D3D11_USAGE_DEFAULT;" << std::endl;
				*classFile << "bufferDesc.BindFlags = D3D11_BIND_UNORDERED_ACCESS | D3D11_BIND_SHADER_RESOURCE;" << std::endl;
				*classFile << "bufferDesc.MiscFlags = D3D11_RESOURCE_MISC_BUFFER_STRUCTURED;" << std::endl;
				*classFile << "bufferDesc.StructureByteStride = sizeof( " << baseType() << ") * " << baseCount() << ";" << std::endl;
				*classFile << "m_pd3dDevice->CreateBuffer( &bufferDesc, NULL, &" << name()
					<< " );" << std::endl;
				
				*classFile << std::endl;
				*classFile << "D3D11_SHADER_RESOURCE_VIEW_DESC srvDesc;" << std::endl;
				*classFile << "memset( &srvDesc, 0, sizeof( srvDesc ) );" << std::endl;
				*classFile << "srvDesc.Format = DXGI_FORMAT_UNKNOWN;" << std::endl;
				*classFile << "srvDesc.ViewDimension = D3D11_SRV_DIMENSION_BUFFER;" << std::endl;
				*classFile << "srvDesc.Buffer.ElementWidth = " << sizeString() << ";" << std::endl;
				*classFile << "srvDesc.Buffer.ElementOffset = 0;" << std::endl;
				*classFile << "m_pd3dDevice->CreateShaderResourceView( " << name()
					<< ", &srvDesc, &" << SRVname() << " );" << std::endl;

				*classFile << std::endl;
				*classFile << "D3D11_UNORDERED_ACCESS_VIEW_DESC uavDesc;" << std::endl;
				*classFile << "memset( &uavDesc, 0, sizeof( uavDesc ) );" << std::endl;
				*classFile << "uavDesc.Format = DXGI_FORMAT_UNKNOWN;" << std::endl;
				*classFile << "uavDesc.ViewDimension = D3D11_UAV_DIMENSION_BUFFER;" << std::endl;
				*classFile << "uavDesc.Buffer.NumElements = " << sizeString() << ";" << std::endl;
				*classFile << "m_pd3dDevice->CreateUnorderedAccessView( " << name() 
					<< ", &uavDesc, &" << UAVname() << " );" << std::endl;


				*classFile << "}" << std::endl;
			}
		}
	}


	void
		Buffer::selfLoadFromDirectComputeDevice( ofstream * classFile )
	{
		*classFile << baseType() << " * _" << name() << " = new " << baseType()
			<< "[ ( " << totalSize() << " ) * "
			<< baseCount() << " ];" << std::endl;
		*classFile << "m_pd3dImmediateContext->CopyResource( " 
			<< name() << "Mappable, " << name() << " );" << std::endl;
		*classFile << "m_pd3dImmediateContext->Map( "
			<< name() << "Mappable, 0, D3D11_MAP_READ, 0, &MappedResource );" << std::endl;
		*classFile << "memcpy( _" << name() << ", MappedResource.pData, "
			//<< this->byteWidthString() << " );" << std::endl;
			<< totalSizeT() << " );" << std::endl;
		*classFile << "m_pd3dImmediateContext->Unmap( "
			<< name() << "Mappable, 0 );" << std::endl;
	}



	void
		Buffer::selfLoadFromOpenCLDevice( ofstream * classFile )
	{
		*classFile << baseType() << " * _" << name() << " = new " << baseType()
			<< "[ (int)( " << totalSize() << " ) * "
			<< baseCount() << " ];" << std::endl;
		if( !isImage() ){
			*classFile << "queue.enqueueReadBuffer( " << name() << ", CL_TRUE, 0, "
				<< totalSizeT() << ", _" << name() << " );" << std::endl;
		}
		else{
			*classFile << "{" << std::endl;
			*classFile << "#if 0//OpenCL bug workaround" << std::endl;
			*classFile << "cl::size_t< 3 > offset( 0, 0, 0 );" << std::endl;
			*classFile << "cl::size_t< 3 > region( ";
			*classFile << dimensionsSizeT();
			if( 2 == dimensionality() ){
				*classFile << ", 1";
			}
			*classFile << " );" << std::endl;
			*classFile << "#else" << std::endl;
			*classFile << "cl::size_t< 3 > offset; offset[ 0 ] = offset[ 1 ] = offset[ 2 ] = 0;" << std::endl;
			*classFile << "cl::size_t< 3 > region;" << std::endl;

			ASTExpr v;
			v = sizeX();
			*classFile << "region[ 0 ] = ( " << stringify( v ) << " );" << std::endl;
			v = sizeY();
			*classFile << "region[ 1 ] = ( " << stringify( v )  << " );" << std::endl;
			v = sizeZ();
			if( 2 == dimensionality() ){
				*classFile << "region[ 2 ] = 1;" << std::endl;
			}
			else{
				*classFile << "region[ 2 ] = ( " << stringify( v )  << " );" << std::endl;
			}
			*classFile << "#endif" << std::endl;
			*classFile << "queue.enqueueReadImage( " << name() << ", CL_TRUE, offset, region, 0, 0, ";
			*classFile << "_" << name() << ", ";
			*classFile << "NULL, NULL );" << std::endl;
			*classFile << "}" << std::endl;
		}
		*classFile << "queue.finish();" << std::endl;
	}


	void
		Buffer::selfLoadFromDevice( ofstream * classFile )
	{
		if( addressSpace() == AST_GLOBAL ){
			if( pipeline()->targetIsOpenCL() ){
				selfLoadFromOpenCLDevice( classFile );
			}
			else if( pipeline()->targetIsDirectCompute() ){
				selfLoadFromDirectComputeDevice( classFile );
			}
		}
	}



	bool
		Buffer::isUsedAsOutput( Stage * s )
		{
#if 0
			bool result = ( s->outputs().find( name() ) != s->outputs().end() );
			return result;
#else
			NodeList::iterator nIt = s->arguments().begin();
			for( KernelParameterList::iterator kplIt = s->kernel()->kernelParameters().begin();
				kplIt != s->kernel()->kernelParameters().end(); ++kplIt, ++nIt ){
					Node * node = *nIt;
					Buffer * B = dynamic_cast<Buffer *>( node );
					if( B != NULL ){
						if( B->name().compare( name() ) == 0 ){
							return ( (*kplIt).ioClass.compare( "input" ) != 0 );
						}
					}
			}//for
			return false;	
#endif
		}

#if 1
	bool
		Buffer::isUsedAsOutput( Kernel * k )
		{
#if 0
			bool result = ( k->outputs().find( name() ) != k->outputs().end() );
			return result;
#else
			for( KernelParameterList::iterator kplIt = k->kernelParameters().begin();
				kplIt != k->kernelParameters().end(); ++kplIt ){
					if( (*kplIt).name.compare( name() ) == 0 ){
						return( (*kplIt).ioClass.compare( "input" ) != 0 );
					}
			}//for
			return false;
#endif
		}
#endif


	void
		Buffer::setDirectComputeAccessView( 
		ofstream * classFile, 
		std::string lastArg, 
		Stage * s,
		unsigned int & srvNumber,
		unsigned int & uavNumber
		)
	{
		if( addressSpace() != AST_LOCAL ){
			std::string uavName;
			std::string srvName;
			if( 0 == lastArg.compare( "NULL" )){
				uavName = "pUAViewNULL";
				srvName = "pSRViewNULL";
			}
			else{
				uavName = UAVname();
				srvName = SRVname();
			}
			if( isUsedAsOutput( s ) ){
				*classFile << "m_pd3dImmediateContext->CSSetUnorderedAccessViews( "
					<< uavNumber++ << ", 1, &"
					<< uavName << ", " << lastArg << ");" << std::endl;
			}
			else{
				*classFile << "m_pd3dImmediateContext->CSSetShaderResources( "
					<< srvNumber++ << ", 1, &"
					<< srvName << ");" << std::endl;
			}
		}
	}

	void
		Buffer::selfCaptureData( ofstream * classFile, Stage * stage )
	{
		if( addressSpace() == AST_GLOBAL ){
			selfLoadFromDevice( classFile );
			Kernel * k = stage->kernel();
			std::string kernelName;
			if( k != NULL ){
				kernelName = k->name();
			}
			else{
				kernelName = stage->name();
			}
			*classFile << "captureBinaryData( \"" << name() << "\", \""
				<< kernelName << "\", _" << name()
				<< ", " << totalSizeT() << " );" << std::endl;

			*classFile << captureFunction() << "( \"" << name() << "\", \""
				<< kernelName << "\", _" << name()
				<< ", " << totalSize() << " );" << std::endl;

			*classFile << "delete [] _" << name() << ";" << std::endl;
		}
	}

	void
		Buffer::selfTerminate( ofstream * classFile )
	{
		if( addressSpace() == AST_GLOBAL ){
			*classFile << "delete " << name() << ";" << std::endl;
		}
	}

	void
		Buffer::check()
	{
		pipeline()->checkName( name() );
		std::string description = "size expression for buffer " + name();
		pipeline()->checkExpression( description, size() );
	}


	// Class Uniform

	void Uniform::selfDeclare( ofstream * classFile )
	{
		if( type().compare( "string" ) ){
			*classFile << type() << " " << name() << ";" << std::endl;
		}
	}

	void
		Uniform::selfUpdateStruct( ofstream * classFile, Stage * s )
	{
		if( isDefined( "_" )){
			ASTExpr rhs = expression( "_" );
			*classFile << "UniformBuffer" << s->kernel()->name() << "_." << s->kernelUniform( this )
				<< " = " << stringify( rhs ) << ";" << std::endl;
		}
		else{
			*classFile << "UniformBuffer" << s->kernel()->name() << "_." << s->kernelUniform( this )
				<< " = " << name() << ";" << std::endl;
		}
	}

	void
		Uniform::selfDeclareAccessorMutator( ofstream * classFile )
	{
		*classFile << std::endl;
		*classFile << type() << std::endl;
		*classFile << "get" << name() << "( void ){" << std::endl;
		*classFile << "return " << name() << ";" << std::endl;
		*classFile << "}" << std::endl;

		*classFile << std::endl;
		*classFile << "void" << std::endl;
		*classFile << "set" << name() << "( " << type() << " value ){" << std::endl;
		*classFile << name() << " = value;" << std::endl;

		*classFile << "}" << std::endl;
	}


	void
		Uniform::generateDependantEvaluations( ofstream * classFile )
	{
		// be sure to invoke every equation that depends on this one
		EquationMap::iterator eqIt;
		eqIt = pipeline()->equations().begin();
		while( eqIt != pipeline()->equations().end() ){
			Equation * e = eqIt->second;
			ASTExpr v = e->value();
			NameList precedents = expr2NameList( v );
			bool isDependant = false;
			NameList::iterator nIt;
			nIt = precedents.begin();
			while( nIt != precedents.end() ){
				std::string precedentName = *nIt;
				if( 0 == precedentName.compare( name() ) ){
					isDependant = true;
					break;
				}
				nIt++;
			}//while
			if( isDependant ){
				*classFile << "evaluate" << e->name() << "();" << std::endl;
			}
			eqIt++;
		}//while
	}

	void
		Uniform::check()
	{
		pipeline()->checkName( name() );
		if( isDefined( "_" )){
			Maybe<ASTExpr>::MaybeType me = maybeExpr( "_" );
			ASTExpr v = Maybe<ASTExpr>::toJust( me );
			pipeline()->checkExpression( name(), v );
		}
	}

	// Class Image2D

	// Class Image3D

	// Class Kernel

	void
		Kernel::selfDeclare( ofstream * classFile )
	{
		if( pipeline()->targetIsOpenCL() ){
			*classFile << "cl::Kernel " << name() << ";" << std::endl;
		}
		else if( pipeline()->targetIsDirectCompute() ){
			*classFile << "ID3D11ComputeShader * " << name() << ";" << std::endl;
		}
	}

	void
		Kernel::selfInitialize( ofstream * classFile )
	{
		if( pipeline()->targetIsOpenCL() ){
			*classFile << name() << " = cl::Kernel( program, \""
				<< name() << "\" );" << std::endl;
		}
		else if( pipeline()->targetIsDirectCompute() ){
			if( isDefined( "sourceFunction" )){
				Maybe<ASTExpr>::MaybeType me = maybeExpr( "sourceFunction" );
				ASTExpr v = Maybe<ASTExpr>::toJust( me );
				std::string sourceStringText = stringify( v );
				*classFile << "createComputeShader ( \"" << name() << "\", &"
					<< name() << ", " << sourceStringText << "(), false );" << std::endl;
			}
			else {
				*classFile << "createComputeShader( \"" << name() << "\", &" 
					<< name() << ", \"" << name() << ".hlsl\", true );" << std::endl;
			}
		}
	}
	

	void
		Kernel::selfGenerateOpenCLHeader( ofstream * kernelFile )
	{
		*kernelFile << std::endl;
		*kernelFile << "__kernel void " << name() << "(" << std::endl;
		bool first = true;
		for( KernelParameterList::iterator kplIt = kernelParameters().begin();
			kplIt != kernelParameters().end(); kplIt++ ){
				if( !first ){
					*kernelFile << "," << std::endl;
				}
				first = false;
				KernelParameter kernelParameter= *kplIt;

				switch( kernelParameter.addressSpace ) {
		case AST_GLOBAL:		
			*kernelFile << "__global ";
			break;
		case AST_LOCAL:
			*kernelFile << "__local ";
			break;
		case AST_UNIFORM:
			break;
		default:
			break;
				}//switch

				bool isBuffer = ( NULL != dynamic_cast<Buffer *>( kernelParameter.node ));
				*kernelFile << kernelParameter.type 
					<< ( isBuffer ? " * " : " " )
					<< kernelParameter.name;
		}//for

#if 0//old
		// input buffers
		BufferMap::iterator bufIt;
		bufIt = inputs().begin();
		bool hasInputBuffers = false;
		while( bufIt != inputs().end() ){
			if( hasInputBuffers ){
				*kernelFile << "," << std::endl;
			}
			hasInputBuffers = true;
			Buffer * B  = bufIt->second;
			*kernelFile << B->kernelSignature( this );
			bufIt++;
		}//while

		// input uniforms
		UniformMap::iterator unIt;
		unIt = uniforms().begin();
		bool hasInputUniforms = false;
		while( unIt != uniforms().end() ){
			if( hasInputBuffers || hasInputUniforms ){
				*kernelFile << "," << std::endl;
			}
			hasInputUniforms = true;
			Uniform * U = unIt->second;
			*kernelFile << U->type() << " " << U->name();
			unIt++;
		}//while

		// output buffers
		bufIt = outputs().begin();
		bool hasOutputBuffers = false;
		while( bufIt != outputs().end() ){
			Buffer * B  = bufIt->second;
			if( inputs().find( B->name() ) == inputs().end() ){
				if( hasInputBuffers || hasInputUniforms || hasOutputBuffers ){
					*kernelFile << "," << std::endl;
				}
				hasOutputBuffers = true;
				*kernelFile << B->kernelSignature( this );
			}
			bufIt++;
		}//while
#endif

		*kernelFile << std::endl;
		*kernelFile << ")" << std::endl;
		*kernelFile << "{" << std::endl;
		*kernelFile << "}" << std::endl;
	}


	void
		Kernel::selfGenerateUniformBuffers( ofstream * classFile )
	{
		if( pipeline()->targetIsDirectCompute() ){
			*classFile << std::endl;
			*classFile << "struct " << "{" << std::endl;
			UniformMap::iterator unIt;
#if 0
			unIt = uniforms().begin();
			int totalSize = 0;
			while( unIt != uniforms().end() ){
				Uniform * u = unIt->second;
				u->selfDeclare( classFile );
				totalSize += u->size();
				unIt++;
			}//while
#else
			int totalSize = 0;
			for( KernelParameterList::iterator kplIt = kernelParameters().begin();
				kplIt != kernelParameters().end(); ++kplIt ){
					Uniform * u = dynamic_cast<Uniform *>( (*kplIt).node );
					if( u != NULL ){
						u->selfDeclare( classFile );
						totalSize += u->size();
					}
			}//for
#endif
			int paddedSize = ( ( ( totalSize - 1 ) / 16 ) + 1 ) * 16;
			int padding = paddedSize - totalSize;
			if( padding > 0 ){
				*classFile << "char __padding[ " << padding << " ];" << std::endl;
			}
			*classFile << "} UniformBuffer" << name() << "_;" << std::endl;
			*classFile << std::endl;
		}
	}


	void
		Kernel::selfGenerateDirectComputeHeader( ofstream * kernelFile )
	{
		*kernelFile << std::endl;
		Maybe<ASTExpr>::MaybeType me = maybeExpr( "workgroupSize" );
		ASTExpr v = Maybe<ASTExpr>::toJust( me );
		std::string workgroupSizeText = stringify( v );
		*kernelFile << "[numthreads(" << workgroupSizeText << ",1,1)]" << std::endl;//to do handle multidimensional issue
		*kernelFile << "void " << name() 
			<< "( uint3 Gid : SV_GroupID, uint3 DTid : SV_DispatchThreadID, uint3 GTid : SV_GroupThreadID, uint GI : SV_GroupIndex )" 
			<< std::endl;
		*kernelFile << "{" << std::endl;
		*kernelFile << "}" << std::endl;
	}

	void
		Kernel::selfGenerateHeader( ofstream * kernelFile )
	{
		if( pipeline()->targetIsOpenCL() ){
			selfGenerateOpenCLHeader( kernelFile );
		}
		else if( pipeline()->targetIsDirectCompute() ){
			selfGenerateDirectComputeHeader( kernelFile );
		}
	}

	void
		Kernel::check()
	{
		pipeline()->checkName( name() );
#if 0//old
		BufferMap::iterator inIt = inputs().begin();
		while( inIt != inputs().end() ){
			Buffer * B = inIt->second;
			pipeline()->checkName( B->name() );
			inIt++;
		}//while
		UniformMap::iterator unIt = uniforms().begin();
		while( unIt != uniforms().end() ){
			Uniform * U = unIt->second;
			pipeline()->checkName( U->name() );
			unIt++;
		}//while
		BufferMap::iterator outIt = outputs().begin();
		while( outIt != outputs().end() ){
			Buffer * B = outIt->second;
			pipeline()->checkName( B->name() );
			outIt++;
		}//while
#else
		for( KernelParameterList::iterator kplIt = kernelParameters().begin();
			kplIt != kernelParameters().end(); ++kplIt ){
			KernelParameter kernelParameter = *kplIt;
			pipeline()->checkName( kernelParameter.name );
		}//for
#endif
	}

	// Class StageBaseClass

#if 0
	void
		StageBaseClass::importMapsFromNameList( 
		NameList nameList,
		BufferMap & buffers,
		UniformMap & uniforms,
		BufferVector &buffersVector,
		UniformVector &uniformsVector,
		bool isInput
		)
	{
		NameList::iterator nameIt;
		nameIt = nameList.begin();
		while( nameIt != nameList.end() ){
			std::string name = *nameIt;
			Pipeline * p = pipeline();
			BufferMap bm = p->buffers();
			BufferMap::iterator bit = bm.find( name );
			if( pipeline()->buffers().find( name ) != pipeline()->buffers().end() ){
				Buffer * B = pipeline()->buffers()[ name ];
				buffers[ name ] = B;
				buffersVector.push_back( B );
				if( isInput ){
					B->isInputIs( true );
				}
				else{
					B->isOutputIs( true );
				}
			}
			else if( pipeline()->image2Ds().find( name ) != pipeline()->image2Ds().end() ){
				Image2D * I = pipeline()->image2Ds()[ name ];
				buffers[ name ] = I;
				buffersVector.push_back( I );
				if( isInput ){
					I->isInputIs( true );
				}
				else{
					I->isOutputIs( true );
				}
			}
			else if( pipeline()->image3Ds().find( name ) != pipeline()->image3Ds().end() ){
				Image3D * I = pipeline()->image3Ds()[ name ];
				buffers[ name ] = I;
				buffersVector.push_back( I );
				if( isInput ){
					I->isInputIs( true );
				}
				else{
					I->isOutputIs( true );
				}
			}
			else if( pipeline()->uniforms().find( name ) != pipeline()->uniforms().end() ){
				Uniform * U = pipeline()->uniforms()[ name ];
				uniforms[ name ] = U;
				uniformsVector.push_back( U );
			}
			else{
				buffers[ name ] = NULL;
				buffersVector.push_back( NULL );
			}
			nameIt++;
		}//while
	}
#endif


	void
		StageBaseClass::importArgumentsFromNameList( 
		NameList nameList,
		NodeList & arguments
		)
	{
		NameList::iterator nameIt;
		nameIt = nameList.begin();
		while( nameIt != nameList.end() ){
			std::string name = *nameIt;
			Pipeline * p = pipeline();
			BufferMap bm = p->buffers();
			BufferMap::iterator bit = bm.find( name );
			if( pipeline()->buffers().find( name ) != pipeline()->buffers().end() ){
				Node * B = pipeline()->buffers()[ name ];
				arguments.push_back( B );
			}
			else if( pipeline()->image2Ds().find( name ) != pipeline()->image2Ds().end() ){
				Node * I = pipeline()->image2Ds()[ name ];
				arguments.push_back( I );
			}
			else if( pipeline()->image3Ds().find( name ) != pipeline()->image3Ds().end() ){
				Node * I = pipeline()->image3Ds()[ name ];
				arguments.push_back( I );
			}
			else if( pipeline()->uniforms().find( name ) != pipeline()->uniforms().end() ){
				Node * U = pipeline()->uniforms()[ name ];
				arguments.push_back( U );
			}
			else{
				arguments.push_back( NULL );
			}
			nameIt++;
		}//while
	}

#if 0//old
	void
		StageBaseClass::importMapsFromASTTypeVarList( 
		ASTTypeVarList parameterList,
		BufferMap & buffers,
		UniformMap & uniforms,
		BufferVector &buffersVector,
		UniformVector &uniformsVector,
		bool isInput
		)
	{
		ASTTypeVarList::iterator avtIt;
		avtIt = parameterList.begin();
		while( avtIt != parameterList.end() ){
			ASTTypeVarElement tuple = *avtIt;
			ADDRESS_SPACE addressSpace = boost::get<0>( tuple );
			std::string type = boost::get<1>( tuple );
			std::string var = boost::get<2>( tuple );
			int tokenType = boost::get< 3 >( tuple );

			if( tokenType == Parser::UNIFORM ){
				ASTAttributeList a;
				Uniform * U;
				U = new Uniform( var, type, a, pipeline() );
				uniforms[ var ] = U;
				uniformsVector.push_back( U );
			}
			else if( tokenType == Parser::BUFFER ){
				ASTAttributeList a;
				Buffer * B;
				B = new Buffer( var, addressSpace, type, ASTExpr( 0 ), a, pipeline() );

				buffers[ var ] = B;
				buffersVector.push_back( B );
				if( isInput ){
					B->isInputIs( true );
				}
				else{
					B->isOutputIs( true );
				}
			}
			else if( tokenType == Parser::IMAGE2D ){
				ASTAttributeList a;
				Image2D * I;
				I = new Image2D( var, addressSpace, type, ASTExpr( 0 ), ASTExpr( 0 ), a, pipeline() );

				buffers[ var ] = I;
				buffersVector.push_back( I );
				if( isInput ){
					I->isInputIs( true );
				}
				else{
					I->isOutputIs( true );
				}
			}
			else if( tokenType == Parser::IMAGE3D ){
				ASTAttributeList a;
				Image3D * I;
				I = new Image3D( var, addressSpace, type, ASTExpr( 0 ), ASTExpr( 0 ), ASTExpr( 0 ), 
					a, pipeline() );

				buffers[ var ] = I;
				buffersVector.push_back( I );
				if( isInput ){
					I->isInputIs( true );
				}
				else{
					I->isOutputIs( true );
				}
			}
			avtIt++;
		}//while
	}
#endif

	void
		StageBaseClass::importParametersFromASTTypeVarList( 
		ASTTypeVarList parameterList,
		KernelParameterList & kernelParameters
		)
	{
		ASTTypeVarList::iterator avtIt;
		avtIt = parameterList.begin();
		while( avtIt != parameterList.end() ){
			ASTTypeVarElement tuple = *avtIt;
			ADDRESS_SPACE addressSpace = boost::get<0>( tuple );
			std::string type = boost::get<1>( tuple );
			std::string var = boost::get<2>( tuple );
			int tokenType = boost::get< 3 >( tuple );
			std::string ioClass = boost::get<4>( tuple );

			KernelParameter kernelParameter;
			kernelParameter.addressSpace = addressSpace;
			kernelParameter.type = type;
			kernelParameter.tokenType = tokenType;
			kernelParameter.ioClass = ioClass;
			kernelParameter.name = var;

			if( tokenType == Parser::UNIFORM ){
				ASTAttributeList a;
				Uniform * U;
				U = new Uniform( var, type, a, pipeline() );
				kernelParameter.node = U;
			}
			else if( tokenType == Parser::BUFFER ){
				ASTAttributeList a;
				Buffer * B;
				B = new Buffer( var, addressSpace, type, ASTExpr( 0 ), a, pipeline() );
				kernelParameter.node = B;
			}
			else if( tokenType == Parser::IMAGE2D ){
				ASTAttributeList a;
				Image2D * I;
				I = new Image2D( var, addressSpace, type, ASTExpr( 0 ), ASTExpr( 0 ), a, pipeline() );
				kernelParameter.node = I;
			}
			else if( tokenType == Parser::IMAGE3D ){
				ASTAttributeList a;
				Image3D * I;
				I = new Image3D( var, addressSpace, type, ASTExpr( 0 ), ASTExpr( 0 ), ASTExpr( 0 ), 
					a, pipeline() );
				kernelParameter.node = I;
			}
			kernelParameters.push_back( kernelParameter );
			avtIt++;
		}//while
	}


	// Class Stage

	Stage::Stage( 
		std::string name,
		std::string kernelName,
#if 0
		NameList inputNames, 
		NameList outputNames, 
#else
		NameList argumentList,
#endif
		ASTAttributeList attributes ,
		Pipeline * pipeline
		)
	{
		pipelineIs( pipeline );
		name_ = name;
		//kernel_ = pipeline->kernels()[ kernelName ];
		kernel_ = NULL;
		if( pipeline->kernels().find( kernelName ) != pipeline->kernels().end() ){
			kernel_ = pipeline->kernels()[ kernelName ];
		}
#if 0
		importMapsFromNameList( inputNames, inputs(), uniforms(), inputsVector(), uniformsVector(), true );
		importMapsFromNameList( outputNames, outputs(), uniforms(), outputsVector(), uniformsVector(), false );
#else
		importArgumentsFromNameList( argumentList, arguments() );
#endif
		importAttributes( attributes );
#if 0
		inputNamesIs( inputNames );
#endif
	}


	void
		Stage::selfDeclareExternal( ofstream * classFile )
	{
		if( isTrue( "external" )){
			*classFile << "unsigned int _Stage_" << name() << "( QueueType queue );" << std::endl;
		}
		if( isTrue( "branchSelector" )){
			*classFile << "unsigned int _Stage_" << name() << "( QueueType queue );" << std::endl;
		}
	}

	void
		Stage::selfInvokeExternal( ofstream * classFile )
	{
		if( isTrue( "external" )){
			*classFile << "unsigned int _run" << name() << "( __SIGNATURE__ )" << std::endl;
			*classFile << "{" << std::endl;

			pipeline()->implementQueueFinish( classFile );
			pipeline()->implementQueueRestart( classFile );

			*classFile << "_Stage_" << name() << "( queue );" << std::endl;
			*classFile << "return 0;" << std::endl;
			*classFile << "}" << std::endl;
		}
	}


	void
		Stage::selfImplementExternal( ofstream * extFile )
	{
		if( isTrue( "external" )){
			*extFile << std::endl;
			*extFile << "unsigned int " << pipeline()->name() << "::_Stage_" << name() << "( QueueType queue ){" << std::endl;

			*extFile << "// Please edit this to implement stage " << name() << std::endl;
			*extFile << "return 0;" << std::endl;
			*extFile << "}" << std::endl;
		}
		else if( isTrue( "branchSelector" )){
			*extFile << std::endl;
			*extFile << "unsigned int " << pipeline()->name() << "::_Stage_" << name() << "( QueueType queue ){" << std::endl;
			*extFile << "// initialize the map of return values" << std::endl;
			*extFile << "static std::map< std::string, unsigned int > returnMap;" << std::endl;
			*extFile << "if( returnMap.size() == 0 ){" << std::endl;
			GraphVertexChildren children = graphVertex()->children();
			std::string exampleName;
			for( unsigned int i = 0; i < children.size(); ++i ){
				Stage * successor = children[ i ]->stage();
				*extFile << "returnMap[ \"" << successor->name() << "\" ] = 1 << " << i << ";" << std::endl;
				if( !i ){
					exampleName = successor->name();
				}
			}//for
			*extFile << "}" << std::endl;
			*extFile << "//" << std::endl;
			*extFile << "// fill in the body of the function here" << std::endl;

			*extFile << "unsigned int result = 0; // follow no branches" << std::endl;
			*extFile << "// return a bit mask to indicate which branches to follow" << std::endl;
			*extFile << "// use a bitwise-or operator | to return multiple branches" << std::endl;
			*extFile << "// for example to pass control to stage " << exampleName << " do this:" << std::endl;
			*extFile << "result = returnMap[ \"" << exampleName << "\" ];" << std::endl;
			*extFile << "return result;" << std::endl;
			*extFile << "}" << std::endl;
		}
	}


	void
		Stage::declareSynchronizationObject( ofstream * classFile )
	{
		*classFile << std::endl;
		*classFile << "SynchronizationObject _sync_" << stringValue( "barrier" ) << ";" << std::endl;
		*classFile << "Mutex _mutex_" << stringValue( "barrier" ) << ";" << std::endl;
		*classFile << "void _initialize_sync_" << stringValue( "barrier" ) << "(){" << std::endl;
		*classFile << "_sync_" << stringValue( "barrier" ) << ".requiredParticipants = 0;" << std::endl;
		*classFile << "_sync_" << stringValue( "barrier" ) << ".currentParticipants = 0;" << std::endl;
		*classFile << "_sync_" << stringValue( "barrier" ) << ".completed = false;" << std::endl;
		*classFile << "#ifdef _WIN32" << std::endl;
		*classFile << "_mutex_" << stringValue( "barrier" ) << " = CreateMutex( NULL, false, (LPCTSTR)\""
			<< stringValue( "barrier" ) << "\" );" << std::endl;
		*classFile << "#else" << std::endl;
		*classFile << "pthread_mutex_init( &_mutex_" << stringValue( "barrier" ) << ", NULL );" << std::endl;
		*classFile << "#endif" << std::endl;
		*classFile << "}" << std::endl;
	}


	void 
		Stage::selfDeclareTest( ofstream * classFile )
	{
		*classFile << "bool testBegin" << name() << "( QueueType queue );" << std::endl;
		*classFile << "bool testEnd" << name() << "( QueueType queue );" << std::endl;
	}


	void
		Stage::selfImplementTest( ofstream * testFile )
	{
		*testFile << std::endl;
		*testFile << "bool" << std::endl;
		*testFile << pipeline()->name() << "::testBegin" << name() << "( QueueType queue )" << std::endl;
		*testFile << "{" << std::endl;

#if 0
		if( inputs().size() ){
			BufferMap::iterator bufIt;
			bufIt = inputs().begin();
			while( bufIt != inputs().end() ){
				Buffer * B = bufIt->second;
				if( B->addressSpace() == AST_GLOBAL ){
					*testFile << "#ifdef TARGET_OPENCL" << std::endl;
					B->selfLoadFromOpenCLDevice( testFile );
					*testFile << "#elif defined TARGET_DIRECTCOMPUTE" << std::endl;
					B->selfLoadFromDirectComputeDevice( testFile );
					*testFile << "#endif" << std::endl;
				}
				bufIt++;
			}//while
		}
#else
		if( arguments().size() ){
			KernelParameterList::iterator kplIt = kernel()->kernelParameters().begin();
			for( NodeList::iterator nIt = arguments().begin();
				nIt != arguments().end(); ++nIt ){
					KernelParameter kernelParameter = *kplIt++;
					if( kernelParameter.ioClass.compare( "output" ) != 0 ){//input or inout
						Node * node = *nIt;
						Buffer * B = dynamic_cast<Buffer *>( node );
						if( B != NULL ){
							if( B->addressSpace() == AST_GLOBAL ){
								*testFile << "#ifdef TARGET_OPENCL" << std::endl;
								B->selfLoadFromOpenCLDevice( testFile );
								*testFile << "#elif defined TARGET_DIRECTCOMPUTE" << std::endl;
								B->selfLoadFromDirectComputeDevice( testFile );
								*testFile << "#endif" << std::endl;
							}
						}
					}
			}//for
		}
#endif

		*testFile << "\n// edit here to test preconditions for stage " << name() << "\n" << std::endl;

#if 0
		if( inputs().size() ){
			BufferMap::iterator bufIt;
			bufIt = inputs().begin();
			while( bufIt != inputs().end() ){
				Buffer * B = bufIt->second;
				if( B->addressSpace() == AST_GLOBAL ){
					*testFile << "delete [] _" << B->name() << ";" << std::endl;
				}
				bufIt++;
			}//while
		}
#else
		if( arguments().size() ){
			KernelParameterList::iterator kplIt = kernel()->kernelParameters().begin();
			for( NodeList::iterator nIt = arguments().begin();
				nIt != arguments().end(); ++nIt ){
					KernelParameter kernelParameter = *kplIt++;
					if( kernelParameter.ioClass.compare( "output" ) != 0 ){//input or inout
						Node * node = *nIt;
						Buffer * B = dynamic_cast<Buffer *>( node );
						if( B != NULL ){
							if( B->addressSpace() == AST_GLOBAL ){
								*testFile << "delete [] _" << B->name() << ";" << std::endl;
							}
						}
					}
			}//for
		}
#endif

		*testFile << "return true;" << std::endl;
		*testFile << "}" << std::endl;

		*testFile << std::endl;
		*testFile << "bool" << std::endl;
		*testFile << pipeline()->name() << "::testEnd" << name() << "( QueueType queue )" << std::endl;
		*testFile << "{" << std::endl;

#if 0
		if( inputs().size() ){
			BufferMap::iterator bufIt;
			bufIt = inputs().begin();
			while( bufIt != inputs().end() ){
				Buffer * B = bufIt->second;
				if( B->addressSpace() == AST_GLOBAL ){
					*testFile << "#ifdef TARGET_OPENCL" << std::endl;
					B->selfLoadFromOpenCLDevice( testFile );
					*testFile << "#elif defined TARGET_DIRECTCOMPUTE" << std::endl;
					B->selfLoadFromDirectComputeDevice( testFile );
					*testFile << "#endif" << std::endl;
				}
				bufIt++;
			}//while
		}
		if( outputs().size() ){
			BufferMap::iterator bufIt;
			bufIt = outputs().begin();
			while( bufIt != outputs().end() ){
				Buffer * B = bufIt->second;
				if( inputs().find( B->name() ) == inputs().end() ){
					if( B->addressSpace() == AST_GLOBAL ){
						*testFile << "#ifdef TARGET_OPENCL" << std::endl;
						B->selfLoadFromOpenCLDevice( testFile );
						*testFile << "#elif defined TARGET_DIRECTCOMPUTE" << std::endl;
						B->selfLoadFromDirectComputeDevice( testFile );
						*testFile << "#endif" << std::endl;
					}
				}
				bufIt++;
			}//while
		}
#else
		if( arguments().size() ){
			KernelParameterList::iterator kplIt = kernel()->kernelParameters().begin();
			for( NodeList::iterator nIt = arguments().begin();
				nIt != arguments().end(); ++nIt ){
					KernelParameter kernelParameter = *kplIt++;
					Node * node = *nIt;
					Buffer * B = dynamic_cast<Buffer *>( node );
					if( B != NULL ){
						if( B->addressSpace() == AST_GLOBAL ){
							*testFile << "#ifdef TARGET_OPENCL" << std::endl;
							B->selfLoadFromOpenCLDevice( testFile );
							*testFile << "#elif defined TARGET_DIRECTCOMPUTE" << std::endl;
							B->selfLoadFromDirectComputeDevice( testFile );
							*testFile << "#endif" << std::endl;
						}
					}
			}//for
		}
#endif

		*testFile << "\n// edit here to test postconditions for stage " << name() << "\n" << std::endl;

#if 0
		if( inputs().size() ){
			BufferMap::iterator bufIt;
			bufIt = inputs().begin();
			while( bufIt != inputs().end() ){
				Buffer * B = bufIt->second;
				if( B->addressSpace() == AST_GLOBAL ){
					*testFile << "delete [] _" << B->name() << ";" << std::endl;
				}
				bufIt++;
			}//while
		}
		if( outputs().size() ){
			BufferMap::iterator bufIt;
			bufIt = outputs().begin();
			while( bufIt != outputs().end() ){
				Buffer * B = bufIt->second;
				if(( inputs().find( B->name() ) == inputs().end() )
					&& ( B->addressSpace() == AST_GLOBAL )){
						*testFile << "delete [] _" << B->name() << ";" << std::endl;
				}
				bufIt++;
			}//while
		}
#else
		if( arguments().size() ){
			KernelParameterList::iterator kplIt = kernel()->kernelParameters().begin();
			for( NodeList::iterator nIt = arguments().begin();
				nIt != arguments().end(); ++nIt ){
					KernelParameter kernelParameter = *kplIt++;
					Node * node = *nIt;
					Buffer * B = dynamic_cast<Buffer *>( node );
					if( B != NULL ){
						if( B->addressSpace() == AST_GLOBAL ){
							*testFile << "delete [] _" << B->name() << ";" << std::endl;
						}
					}
			}//for
		}
#endif

		*testFile << "return true;" << std::endl;
		*testFile << "}" << std::endl;	
	}


	void
		Stage::selfGenerateSetArgs( ofstream * classFile )
	{
		int count  = 0;
		if( !isDeclaredExternally() ){
			*classFile << "// Stage " << name() << std::endl;
			for( NodeList::iterator nIt = arguments().begin();
				nIt != arguments().end(); ++nIt ){
					Node * node = *nIt;
					Buffer * B = dynamic_cast<Buffer *>( node );
					if( B != NULL ){
						if( B->addressSpace() == AST_GLOBAL ){
							*classFile << kernel()->name() << ".setArg( " << count++
								<< ", " << B->name() << " );" << std::endl;
						}
						else{
							*classFile << kernel()->name() << ".setArg( " << count++
								<< ", cl::__local( " << B->totalSize()
								<< "));//" << B->name() << std::endl;
						}
					}
					Uniform * U = dynamic_cast<Uniform *>( node );
					if( U != NULL ){
						U->selfUpdate( classFile );
						*classFile << kernel()->name() << ".setArg( " << count++ << ", "
							<< U->name() << " );" << std::endl;
					}
			}//for

#if 0
			BufferMap::iterator bufIt;
			// input buffers
			bufIt = inputs().begin();
			while( bufIt != inputs().end() ){
				Buffer * B  = bufIt->second;
				if( B->addressSpace() == AST_GLOBAL ){
					*classFile << kernel()->name() << ".setArg( " << count++
						<< ", " << B->name() << " );" << std::endl;
				}
				else{
					*classFile << kernel()->name() << ".setArg( " << count++
						<< ", cl::__local( " << B->totalSize()
						<< "));//" << B->name() << std::endl;
				}
				bufIt++;
			}//while

			// input uniforms
			UniformMap::iterator unIt;
			unIt = uniforms().begin();
			while( unIt != uniforms().end() ){
				Uniform * U = unIt->second;
				U->selfUpdate( classFile );
				*classFile << kernel()->name() << ".setArg( " << count++ << ", "
					<< U->name() << " );" << std::endl;
				unIt++;
			}//while

			// output buffers
			bufIt = outputs().begin();
			while( bufIt != outputs().end() ){
				Buffer * B  = bufIt->second;
				if( inputs().find( B->name() ) == inputs().end() ){
					if( B->addressSpace() == AST_GLOBAL ){
						*classFile << kernel()->name() << ".setArg( " << count++
							<< ", " << B->name() << " );" << std::endl;
					}
					else{
						*classFile << kernel()->name() << ".setArg( " << count++
							<< ", cl::__local( " << B->totalSize()
							<< "));//" << B->name() << std::endl;
					}
				}
				bufIt++;
			}//while

#endif
		}
	}


	void
		Stage::selfTerminate( ofstream * classFile )
	{
		if( !isDeclaredExternally() ){
			*classFile << "delete " << kernel()->name() << ";" << std::endl;
		}
	}


	void
		Stage::selfCaptureInputs( ofstream * classFile )
	{
#if 1
		if( kernel() != NULL ){
			*classFile << "{" << std::endl;
			KernelParameterList::iterator kplIt = kernel()->kernelParameters().begin();
			for( NodeList::iterator nIt = arguments().begin();
				nIt != arguments().end(); ++nIt, kplIt++ ){
					KernelParameter kernelParameter = *kplIt;
					if( kernelParameter.ioClass.compare( "output" )){//input, inout
						Buffer * B = dynamic_cast<Buffer *>( *nIt );
						if( B != NULL ){
							B->selfCaptureData( classFile, this );
						}
					}
			}//for
			*classFile << "}" << std::endl;
		}
#else//old
		BufferMap::iterator bufIt;
		bufIt = inputs().begin();
		while( bufIt != inputs().end() ){
			Buffer * B = bufIt->second;
			B->selfCaptureData( classFile, this );
			bufIt++;
		}//while
#endif
	}


	void
		Stage::selfCaptureOutputs( ofstream * classFile )
	{
#if 1
		if( kernel() != NULL ){
			*classFile << "{" << std::endl;
			KernelParameterList::iterator kplIt = kernel()->kernelParameters().begin();
			for( NodeList::iterator nIt = arguments().begin();
				nIt != arguments().end(); ++nIt, kplIt++ ){
					KernelParameter kernelParameter = *kplIt;
					if( kernelParameter.ioClass.compare( "input" )){//output, inout
						Buffer * B = dynamic_cast<Buffer *>( *nIt );
						if( B != NULL ){
							B->selfCaptureData( classFile, this );
						}
					}
			}//for
			*classFile << "}" << std::endl;
		}
#else//old
		BufferMap::iterator bufIt;
		bufIt = outputs().begin();
		while( bufIt != outputs().end() ){
			Buffer * B = bufIt->second;
			B->selfCaptureData( classFile, this );
			bufIt++;
		}//while
		*classFile << "}" << std::endl;
#endif
	}


	void
		Stage::selfDeclareDirectComputeBuffers( ofstream * classFile )
	{
		if( false == isDeclaredExternally() ){
			*classFile << "ID3D11Buffer * UniformBuffer" << kernel()->name() << ";" << std::endl;
		}
	}


	void
		Stage::selfInitializeUniformBuffers( ofstream * classFile )
	{
		if( false == isDeclaredExternally() ){
			*classFile << "// create Uniform argument buffer for stage " << name() << std::endl;
			*classFile << "D3D11_BUFFER_DESC UniformBufferDesc" << name() << ";" << std::endl;
			*classFile << "memset( &UniformBufferDesc" << name() 
				<< ", 0, sizeof( UniformBufferDesc" << name() << " ));" << std::endl;
			*classFile << "UniformBufferDesc" << name() 
				<< ".ByteWidth = sizeof( UniformBuffer" << kernel()->name() << "_ );" << std::endl;
			*classFile << "UniformBufferDesc" << name() << ".Usage = D3D11_USAGE_DYNAMIC;" << std::endl;
			*classFile << "UniformBufferDesc" << name() << ".BindFlags = D3D11_BIND_CONSTANT_BUFFER;" << std::endl;
			*classFile << "UniformBufferDesc" << name() << ".CPUAccessFlags = D3D11_CPU_ACCESS_WRITE;" << std::endl;
			*classFile << "m_pd3dDevice->CreateBuffer( &UniformBufferDesc" << name() 
				<< ", NULL, &UniformBuffer" << kernel()->name() << ");" << std::endl;
		}
	}


	void
		Stage::selfUpdateUniforms( ofstream * classFile )
	{
		if( false == isDeclaredExternally() ){
			*classFile << "// update Uniform arguments for " << name() << std::endl;
#if 0
			UniformMap::iterator unIt;
			unIt = uniforms().begin();
			while( unIt != uniforms().end() ){
				Uniform * U = unIt->second;
				U->selfUpdate( classFile );
				U->selfUpdateStruct( classFile, this );
				unIt++;
			}//while
#else
			for( NodeList::iterator nIt = arguments().begin();
				nIt != arguments().end(); ++nIt ){
					Uniform * U = dynamic_cast<Uniform *>( *nIt );
					if( U != NULL ){
						U->selfUpdate( classFile );
						U->selfUpdateStruct( classFile, this );
					}
			}//for
#endif
			*classFile << "// map Uniforms for compute shader " << kernel()->name() << std::endl;
			*classFile << "D3D11_MAPPED_SUBRESOURCE mappedResource" << name() << " = { 0 };" << std::endl;
			*classFile << "m_pd3dImmediateContext->Map( UniformBuffer" << kernel()->name ()
				<< ", 0, D3D11_MAP_WRITE_DISCARD, 0, &mappedResource" << name() << " );" << std::endl;
			*classFile << "memcpy( mappedResource" << name() << ".pData, &UniformBuffer" << kernel()->name() 
				<< "_, sizeof( UniformBuffer" << kernel()->name() << "_ ));" << std::endl;
			*classFile << "m_pd3dImmediateContext->Unmap( UniformBuffer" << kernel()->name() << ", 0 );" << std::endl;
			*classFile << "m_pd3dImmediateContext->CSSetConstantBuffers( 0, 1, &UniformBuffer" << kernel()->name()
				<< " );" << std::endl;
		}
	}


	void
		Stage::selfInvokeComputeShader( ofstream * classFile )
	{
		*classFile << "// invoke compute shader " << kernel()->name() << std::endl;
		*classFile << "const unsigned int UAVinitialCounts = 0;" << std::endl;
		unsigned int srvNumber = 0;
		unsigned int uavNumber = 0;
#if 0
		BufferMap::iterator bufIt;
		bufIt = inputs().begin();
		while( bufIt != inputs().end() ){
			Buffer * b = bufIt->second;
			b->setDirectComputeAccessView( classFile, "&UAVinitialCounts", this, srvNumber, uavNumber );
			bufIt++;
		}//while
		bufIt = outputs().begin();
		while( bufIt != outputs().end() ){
			Buffer * b = bufIt->second;
			if( inputs().find( b->name() ) == inputs().end() ){
				b->setDirectComputeAccessView( classFile, "&UAVinitialCounts", this, srvNumber, uavNumber );
			}
			bufIt++;
		}//while
#else
		for( NodeList::iterator nIt = arguments().begin();
			nIt != arguments().end(); ++nIt ){
				Buffer * b = dynamic_cast<Buffer *>( *nIt );
				if( b != NULL ){
					b->setDirectComputeAccessView( classFile, "&UAVinitialCounts", this, srvNumber, uavNumber );
				}
		}//for
#endif

		*classFile << "m_pd3dImmediateContext->CSSetShader( " << kernel()->name() << ", NULL, 0 );" << std::endl;
		Maybe<ASTExpr>::MaybeType rangeExpr = pipeline()->maybeExpr( "range" );
		ASTExpr v = Maybe<ASTExpr>::toJust( rangeExpr );
		std::string rangeString = stringify( v );
		Maybe<ASTExpr>::MaybeType workgroupSizeExpr = pipeline()->maybeExpr( "workgroupSize" );
		v = Maybe<ASTExpr>::toJust( workgroupSizeExpr );
		std::string workgroupSizeString = stringify( v );
		*classFile << "int dispatch = ( " << rangeString << " ) / ( " << workgroupSizeString << " );" << std::endl;
		*classFile << "m_pd3dImmediateContext->Dispatch( dispatch, 1, 1 );" << std::endl;

		*classFile << "ID3D11UnorderedAccessView* pUAViewNULL = NULL;" << std::endl;
		*classFile << "ID3D11ShaderResourceView *pSRViewNULL = NULL;" << std::endl;

		srvNumber = 0;
		uavNumber = 0;
#if 0
		bufIt = inputs().begin();
		while( bufIt != inputs().end() ){
			Buffer * b = bufIt->second;
			b->setDirectComputeAccessView( classFile, "NULL", this, srvNumber, uavNumber );
			bufIt++;
		}//while
		bufIt = outputs().begin();
		while( bufIt != outputs().end() ){
			Buffer * b = bufIt->second;
			if( inputs().find( b->name() ) == inputs().end() ){
				b->setDirectComputeAccessView( classFile, "NULL", this, srvNumber, uavNumber );
			}
			bufIt++;
		}//while
#else
		KernelParameterList::iterator kplIt = kernel()->kernelParameters().begin();
		for( NodeList::iterator nIt = arguments().begin();
			nIt != arguments().end(); ++nIt, ++kplIt ){
				Buffer * b = dynamic_cast<Buffer *>( *nIt );
				if( b != NULL ){
					b->setDirectComputeAccessView( classFile, "NULL", this, srvNumber, uavNumber );
				}
		}//for
#endif
	}


	std::string
		Stage::globalNDRange()
	{
		return stringValue( "globalRange" );
	}


	std::string
		Stage::localNDRange()
	{
		return stringValue( "localRange" );
	}


	void
		Stage::selfInvokeRun( ofstream * classFile )
	{
		if( !isDeclaredExternally() && !isSynchronizationStage() ){
			*classFile << std::endl;
			*classFile << "unsigned int" << std::endl;
			*classFile << "_run" << name() <<"( __SIGNATURE__ ){" << std::endl;

			if( pipeline()->targetIsOpenCL() ){
				selfGenerateSetArgs( classFile );
				*classFile << "queue.enqueueNDRangeKernel(" << std::endl;
				*classFile << kernel()->name() << ", cl::NullRange, ";
				// global
				*classFile << globalNDRange() << "," << std::endl;
				//local
				*classFile << "#if defined( __APPLE__ )" << std::endl;
				*classFile << "cl::NullRange, NULL, ( e == NULL ) ? NULL : *e );" << std::endl;
				*classFile << "#else" << std::endl;
				*classFile << localNDRange() << ", NULL, ( e == NULL ) ? NULL : *e );" << std::endl;
				*classFile << "#endif" << std::endl;
				*classFile << "#if !defined( NDEBUG )\nqueue.finish();\n#endif" << std::endl;
			}
			else if( pipeline()->targetIsDirectCompute() ){
				selfUpdateUniforms( classFile );
				selfInvokeComputeShader( classFile );
			}
			*classFile << "return 0;" << std::endl;
			*classFile << "}" << std::endl;
		}
		else if( isTrue( "branchSelector" )){
			*classFile << std::endl;
			*classFile << "unsigned int" << std::endl;
			*classFile << "_run" << name() << "( __SIGNATURE__ ){" << std::endl;
			
			pipeline()->implementQueueFinish( classFile );

			*classFile << "unsigned int result = _Stage_"
				<< name() << "( queue );" << std::endl;
			*classFile << "return result;" << std::endl;
			*classFile << "}" << std::endl;
		}
		else if( isSynchronizationStage() ){
			*classFile << std::endl;
			*classFile << "unsigned int" << std::endl;
			*classFile << "_run" << name() << "( __SIGNATURE__ ){" << std::endl;
			if( isBarrierStage() ){
				*classFile << "barrierSynchronization( _sync_" << stringValue( "barrier" ) << 
					", _mutex_" << stringValue( "barrier" ) << " );" << std::endl;
			}
			*classFile << "return 0;" << std::endl;
			*classFile << "}" << std::endl;
		}
	}



	void
		Stage::selfInvokeTest( ofstream * classFile )
	{
		*classFile << std::endl;
		*classFile << "#if !defined( NDEBUG )" << std::endl;
		*classFile << std::endl;
		*classFile << "unsigned int" << std::endl;
		*classFile << "_test" << name() << "( __SIGNATURE__ ){" << std::endl;
		
		if( !isDeclaredExternally() ){
			pipeline()->implementQueueFinish( classFile );
		}
		
		*classFile << "// Invoke stage " << name() << std::endl;
		*classFile << "bool thisTestPassed = true;" << std::endl;
		*classFile << "std::cerr << \"Testing stage " << name() << ":\" << std::endl;" << std::endl;
		*classFile << "thisTestPassed &= testBegin" << name() << "( queue );" << std::endl;

		if( !isDeclaredExternally() ){
			*classFile << "unsigned int result = _run" << name() << "( queue, e, success );" << std::endl;
			if( pipeline()->targetIsOpenCL() ){
			  pipeline()->implementQueueFinish( classFile );
			}
		}
		else{
			*classFile << "unsigned int result = _run" << name() << "( queue, e, success );" << std::endl;
		}

		if( !isDeclaredExternally() ){
			pipeline()->implementQueueRestart( classFile );
		}

		*classFile << "thisTestPassed &= testEnd" << name() << "( queue );" << std::endl;
		*classFile << "*success &= thisTestPassed;" << std::endl;
		*classFile << "if( !thisTestPassed ){" << std::endl;
		*classFile << "std::cerr << \"Test FAILED in stage " << name() << "\" << std::endl;" << std::endl;
		*classFile << "}else{" << std::endl;
		*classFile << "std::cerr << \"Test SUCCEEDED in stage " << name() << "\" << std::endl;" << std::endl;
		*classFile << "}" << std::endl;
		*classFile << "return result;" << std::endl;
		*classFile << "}" << std::endl;
		*classFile << std::endl;
		*classFile << "#endif" << std::endl;
	}


	void
		Stage::selfInvokeCapture( ofstream * classFile )
	{
		*classFile << std::endl;
		*classFile << "#if !defined( NDEBUG )" << std::endl;
		*classFile << std::endl;
		*classFile << "unsigned int" << std::endl;
		*classFile << "_capture" << name() << "( __SIGNATURE__ ){" << std::endl;

		if( !isDeclaredExternally() ){
			pipeline()->implementQueueFinish( classFile );
		}

		selfCaptureInputs( classFile );

		if( !isDeclaredExternally() ){
			*classFile << "unsigned int result = _run" << name() << "( queue, e, success );" << std::endl;

			if( pipeline()->targetIsOpenCL() ){
				pipeline()->implementQueueFinish( classFile );
			}

		}
		else{
			*classFile << "unsigned int result = _run" << name() << "( queue, e, success );" << std::endl;
		}

		if( !isDeclaredExternally() ){
			pipeline()->implementQueueRestart( classFile );
		}

		selfCaptureOutputs( classFile );

		*classFile << "return result;" << std::endl;
		*classFile << "}" << std::endl;
		*classFile << std::endl;
		*classFile << "#endif" << std::endl;
	}

	void
		Stage::check()
	{
		pipeline()->checkName( name() );
		if( !isDeclaredExternally() && !isSynchronizationStage() ){
			Kernel * K = kernel();
			if( NULL == K ){
				throw std::runtime_error(
					"stage " + name() + " invokes unknown kernel " );
			}
			if( K->kernelParameters().size() != arguments().size() ){
#if 0
				std::cout << "kernel " << K->name() << " has " << K->kernelParameters().size() << " parameters" << std::endl;
				for( KernelParameterList::iterator kplIt = K->kernelParameters().begin();
					kplIt != K->kernelParameters().end(); ++kplIt ){
						std::cout << (*kplIt).name << std::endl;
				}//for
				std::cout << "stage " << name() << " has " << arguments().size() << " arguments" << std::endl;
				for( NodeList::iterator nIt = arguments().begin(); nIt != arguments().end(); ++nIt ){
					Node * node = *nIt;
					Buffer * B = dynamic_cast<Buffer *>( node );
					Uniform * U = dynamic_cast<Uniform *>( node );
					if( B != NULL ){
						std::cout << B->name() << std::endl;
					}
					else if( U != NULL ){
						std::cout << U->name() << std::endl;
					}
				}//for
#endif
				throw std::runtime_error(
					"stage " + name() + " does not agree with kernel "
					+ K->name() + " argument count" );  
			}
			KernelParameterList::iterator kplIt = kernel()->kernelParameters().begin();
			for( NodeList::iterator nIt = arguments().begin();
				nIt != arguments().end(); nIt++, kplIt++ ){
				Node * node = *nIt;
				if( NULL == node ){
					throw std::runtime_error(
						"stage " + name() + " invokes unknown argument for kernel parameter "
						+ kplIt->name );
				}
				// to do check storage class and dataType of argument
			}//for
#if 0
			if( K->inputsVector().size() != inputsVector().size() 
				|| K->uniformsVector().size() != uniformsVector().size()
				|| K->outputsVector().size() != outputsVector().size() ){
					throw std::runtime_error(
						"stage " + name() + " does not agree with kernel "
						+ K->name() + " argument count" );
			}
			// to do - warn about type differences between kernel parameters and stage arguments
			BufferMap::iterator bufIt = inputs().begin();
			while( bufIt != inputs().end() ){
				Buffer * B = bufIt->second;
				if( NULL == B ){
					std::string bufferName = bufIt->first;
					throw std::runtime_error(
						"unknown argument " + bufferName
						+ " to stage " + name() );
				}
				bufIt++;
			}//while
			bufIt = outputs().begin();
			while( bufIt != outputs().end() ){
				Buffer * B = bufIt->second;
				if( NULL == B ){
					std::string bufferName = bufIt->first;
					throw std::runtime_error(
						"unknown argument " + bufferName
						+ " to stage " + name() );
				}
				bufIt++;
			}//while
#endif
		}
	}


	// Class GraphVertex

	void
		GraphVertex::generateVisualization( ofstream * dotFile )
	{
		for( unsigned int i = 0; i < children().size(); ++i ){
			GraphVertex * child = children()[ i ];
			*dotFile << stage()->name() << " -> " << child->stage()->name() << ";" << std::endl;
		}//for

		if( stage()->kernel() != NULL ){
			KernelParameterList::iterator kplIt = stage()->kernel()->kernelParameters().begin();
			for( NodeList::iterator nIt = stage()->arguments().begin();
				nIt != stage()->arguments().end(); nIt++, kplIt++ ){
					Node * node = *nIt;
					Uniform * U = dynamic_cast<Uniform *>( node );
					if( U != NULL ){
						*dotFile << stage()->name() << U->name() << " [label=" << U->name() << ",shape=plaintext];" << std::endl;
						*dotFile << stage()->name() << U->name() << " -> " << stage()->name() << " [style=dotted];" << std::endl;	
					}
					Buffer * B = dynamic_cast<Buffer *>( node );
					if( B != NULL ){
						KernelParameter kernelParameter = *kplIt;
						if( kernelParameter.ioClass.compare( "output" )){//input or inout
							*dotFile << B->name() << " -> " << stage()->name() << " [style=dotted];" << std::endl;
							if( B->addressSpace() == AST_LOCAL ){
								*dotFile << stage()->name() << " -> " << B->name() << " [style=dotted];" << std::endl;
							}
						}
						else if( kernelParameter.ioClass.compare( "input" )){//output or inout
							*dotFile << stage()->name() << " -> " << B->name() << " [style=dotted];" << std::endl;
						}
					}
			}//for
		}

#if 0
		UniformMap::iterator uIt;
		uIt = stage()->uniforms().begin();
		while( uIt != stage()->uniforms().end() ){
			Uniform * U = uIt->second;
			*dotFile << stage()->name() << U->name() << " [label=" << U->name() << ",shape=plaintext];" << std::endl;
			*dotFile << stage()->name() << U->name() << " -> " << stage()->name() << " [style=dotted];" << std::endl;
			uIt++;
		}//while

		BufferMap::iterator bufIt;
		bufIt = stage()->inputs().begin();
		while( bufIt != stage()->inputs().end() ){
			Buffer * B = bufIt->second;
			*dotFile << B->name() << " -> " << stage()->name() << " [style=dotted];" << std::endl;
			if( B->addressSpace() == AST_LOCAL ){
				*dotFile << stage()->name() << " -> " << B->name() << " [style=dotted];" << std::endl;
			}
			bufIt++;
		}//while

		bufIt = stage()->outputs().begin();
		while( bufIt != stage()->outputs().end() ){
			Buffer * B = bufIt->second;
			*dotFile << stage()->name() << " -> " << B->name() << " [style=dotted];" << std::endl;
			bufIt++;
		}//while

#endif
	}

	// Class Graph

	Graph::Graph( std::string name,
		ASTGraphEdgeList edges, 
		ASTAttributeList attributes,
		Pipeline * pipeline
		)
	{
		pipelineIs( pipeline );
		name_ = name;
		edges_ = edges;
		importEdgeList( edges );
		importAttributes( attributes );

	}


	void 
		Graph::importEdgeList( ASTGraphEdgeList edges )
	{
		ASTGraphEdgeList::iterator edgeIt;
		edgeIt = edges.begin();
		while( edgeIt != edges.end() ){
			ASTGraphEdge e = *edgeIt;
			std::string stageName_i = boost::get< 0 >( e );
			std::string stageName_j = boost::get< 1 >( e );
			GraphVertex * vertex_i = vertices_[ stageName_i ];
			if( NULL == vertex_i ){
				Stage * S = pipeline()->stages()[ stageName_i ];
				if( NULL == S ){
					throw std::runtime_error( "graph references unknown stage " + stageName_i );
				}
				vertex_i = new GraphVertex( S, this );
				vertices()[ stageName_i ] = vertex_i;
				S->graphVertexIs( vertex_i );
			}
			GraphVertex * vertex_j= vertices_[ stageName_j ];
			if( NULL == vertex_j ){
				Stage * S = pipeline()->stages()[ stageName_j ];
				if( NULL == S ){
					throw std::runtime_error( "graph references unknown stage " + stageName_i );
				}
				vertex_j = new GraphVertex( S, this );
				vertices()[ stageName_j ] = vertex_j;
				S->graphVertexIs( vertex_j );
			}
			vertex_i->children().push_back( vertex_j );
			edgeIt++;
		}//while
	}


	void
		Graph::generateVisualization( ofstream * dotFile )
	{
		*dotFile << std::endl;
		*dotFile << "digraph " << name() << "{" << std::endl;
		*dotFile << "[rankdir=LR];" << std::endl;

		BufferMap::iterator bufIt;
		bufIt = pipeline()->buffers().begin();
		while( bufIt != pipeline()->buffers().end() ){
			Buffer * B = bufIt->second;
			*dotFile << B->name() << std::endl;
			bufIt++;
		}//while

		StageMap::iterator stageIt;
		stageIt = pipeline()->stages().begin();
		while( stageIt != pipeline()->stages().end() ){
			Stage * S = stageIt->second;
			*dotFile << S->name() << " [shape=box];" << std::endl;
			stageIt++;
		}//while

		GraphVertexMap::iterator gvmIt;
		gvmIt = vertices().begin();
		while( gvmIt != vertices().end() ){
			GraphVertex * GV = gvmIt->second;
			GV->generateVisualization( dotFile );
			gvmIt++;
		}//while

		*dotFile << "}" << std::endl;
	}

	void
		Graph::generateExecute(
		ofstream * classFile,
		std::string prefix
		)
	{
		std::string startName = stringValue( "start" );
		Stage * S = pipeline()->stages().find( startName )->second;
		*classFile << "enqueueGraphBlock( &" << pipeline()->name() <<
			"::_Block_" << prefix << S->name() << ", queue );" << std::endl;
		*classFile << "acquireThread();" << std::endl;
	}


	void
		Graph::generateBasicBlocks( 
		ofstream * classFile, 
		GraphVertex * graphVertex,
		std::string prefix
		)
	{
		if( graphVertex->basicBlocksGenerated() ){
			return;
		}
		graphVertex->basicBlocksGeneratedIs( true );
		Stage * S = graphVertex->stage();
		*classFile << std::endl;
		*classFile << "void _Block" << prefix << S->name() << "( __SIGNATURE__ ){" << std::endl;

		GraphVertex * GV = graphVertex;
		while( 1 == GV->children().size() && !GV->isBranchSelector() ){
			Stage * S = GV->stage();
			*classFile << prefix << S->name() << "( queue, e, success );" << std::endl;
			GV = GV->children()[ 0 ];
		}//while

		if( GV->isBranchSelector() ){
			Stage * SS = GV->stage();
			*classFile << "unsigned int branchSelection = " << prefix << SS->name() << "( queue, e, success );" << std::endl;
			for( unsigned int i = 0; i < GV->children().size(); ++i ){
				GraphVertex *CGV = GV->children()[ i ];
				Stage *CS = CGV->stage();
				//to do invoke this via threads
				*classFile << "if( branchSelection & " << ( 1 << i ) << " ){" << std::endl;
				*classFile << "__PIPELINEMEMBER__ b = &" << pipeline()->name() << "::_Block" << prefix << CS->name() << ";" << std::endl;
				*classFile << "enqueueGraphBlock( b, queue );" << std::endl;
				*classFile << "}" << std::endl;
			}//for
			*classFile << "}// _Block_" << prefix << S->name() << std::endl;
			for( unsigned int i = 0; i < GV->children().size(); ++i ){
				GraphVertex *CGV = GV->children()[ i ];
				Stage *CS = CGV->stage();
				generateBasicBlocks( classFile, CGV, prefix );//recurse
			}//for
		}
		else{
			assert( GV->children().size() == 0 );
			Stage * SS = GV->stage();
			*classFile << prefix << SS->name() << "( queue, e, success );" << std::endl;
			*classFile << "}// _Block_" << prefix << S->name() << std::endl;
		}

	}

	void
		Graph::check()
	{
		pipeline()->checkName( name() );
		GraphVertexMap::iterator gvmIt;
		gvmIt = vertices().begin();
		while( gvmIt != vertices().end() ){
			GraphVertex * graphVertex = gvmIt->second;
			if( NULL == graphVertex->stage() ){
				std::string stageName = gvmIt->first;
				throw std::runtime_error(
					"unknown stage " + stageName
					+ " in graph " + name() );
			}
			gvmIt++;
		}//while
		if( !isDefined( "start" )){
			throw std::runtime_error(
				"no start stage defined for graph " + name() );
		}
		std::string startStageName = stringValue( "start" );
		if( pipeline()->stages().find( startStageName ) == pipeline()->stages().end() ){
			throw std::runtime_error(
				"start stage " + startStageName + " not found for graph " + name() );
		}
	}

	// Pipeline

	void
		Pipeline::implementQueueFinish( ofstream * file )
	{
		if( targetIsOpenCL() ){
			*file << "queue.finish();" << std::endl;
		}
		else if( targetIsDirectCompute() ){
			*file << "if( e != NULL ){" << std::endl;
			*file << "m_pd3dImmediateContext->End( (ID3D11Query *)*e );" << std::endl;
			*file << "while( S_OK != m_pd3dImmediateContext->GetData( (ID3D11Query *)*e, NULL, 0, 0 )){}" << std::endl;
			*file << "ID3D11Query * ep = (ID3D11Query *)*e;" << std::endl;
			*file << "SAFE_RELEASE( ep );" << std::endl;
			*file << "}" << std::endl;
		}
	}

	void
		Pipeline::implementQueueRestart( ofstream * file )
	{
		if( targetIsOpenCL() ){
			// nothing
		}
		else if( targetIsDirectCompute() ){
			*file << "if( e != NULL ){" << std::endl;
			*file << "D3D11_QUERY_DESC queryDesc;" << std::endl;
			*file << "queryDesc.Query = D3D11_QUERY_EVENT;" << std::endl;
			*file << "queryDesc.MiscFlags = 0;" << std::endl;
			*file << "ID3D11Query * eNew;" << std::endl;
			*file << "HRESULT hres = m_pd3dDevice->CreateQuery( &queryDesc, &eNew );" << std::endl;
			*file << "*e = (EventType *)eNew;" << std::endl;
			*file << "}" << std::endl;
		}
	}

	void
		Pipeline::generateCPPCodes( ofstream * classFile, bool isTopLevel )
	{
		bool first = true;
		CodeMap::iterator codeIt;
		codeIt = codes().begin();
		while( codeIt != codes().end() ){
			Code * c = codeIt->second;
			bool showThis = ( c->isDefined( "topLevel" ) == isTopLevel );
			if( c->name().compare( "CPP" ) && showThis ){
				if( first ){
					first = false;
					*classFile << "// User-declared code:" << std::endl;
				}
				*classFile << c->code();
			}
			codeIt++;
		}//while
		if( !first ){
			*classFile << "// End user-declared code" << std::endl;
			*classFile << std::endl;
		}
	}


	void
		Pipeline::generateCLCodes( ofstream * kernelFile )
	{
		bool first = true;
		CodeMap::iterator codeIt;
		codeIt = codes().begin();
		while( codeIt != codes().end() ){
			Code * c = codeIt->second;
			if( c->name().compare( "CL" ) ){
				if( first ){
					first = false;
					*kernelFile << "// User-declared code:" << std::endl;
				}
				*kernelFile << c->code();
			}
			codeIt++;
		}//while
		if( !first ){
			*kernelFile << "// End user-declared code" << std::endl;
			*kernelFile << std::endl;
		}
	}

	void
		Pipeline::getInteropUsage(
		bool & glInterop,
		bool & dxInterop
		)
	{
		// scan the buffers for glInterop or dxInterop true
		BufferMap::iterator bufIt;
		bufIt = buffers().begin();
		glInterop = false;
		dxInterop = false;
		while( bufIt != buffers().end() ){
			Buffer * B = bufIt->second;
			glInterop |= B->isTrue( "glInterop" );
			dxInterop |= B->isTrue( "dxInterop" );
			bufIt++;
		}//while
	}


	void
		Pipeline::generateInteropBools( ofstream * classFile )
	{
		bool glInterop;
		bool dxInterop;
		getInteropUsage( glInterop, dxInterop );

		*classFile << "const bool isGLInterop( void ){" << std::endl;
		*classFile << "#ifndef GL_PIPELINE" << std::endl;
		*classFile << "   return false;" << std::endl;
		*classFile << "#else" << std::endl;
		*classFile << "   return " << ( glInterop ? "true;" : "false;" ) << std::endl;
		*classFile << "#endif" << std::endl;
		*classFile << "}" << std::endl;

		*classFile << "const bool isDXInterop( void ){" << std::endl;
		*classFile << "#ifndef DX_PIPELINE" << std::endl;
		*classFile << "   return false;" << std::endl;
		*classFile << "#else" << std::endl;
		*classFile << "   return " << ( dxInterop ? "true;" : "false;" ) << std::endl;
		*classFile << "#endif" << std::endl;
		*classFile << "}" << std::endl;
	}

	void
		Pipeline::generateBufferDeclarations( ofstream * classFile )
	{
		*classFile << std::endl;
		*classFile << "// Buffers" << std::endl;
		BufferMap::iterator bufIt;
		bufIt = buffers().begin();
		while( bufIt != buffers().end() ){
			Buffer * B = bufIt->second;
			B->selfDeclare( classFile );
			bufIt++;
		}//while
		Image2DMap::iterator i2it;
		i2it = image2Ds().begin();
		while( i2it != image2Ds().end() ){
			Image2D * I = i2it->second;
			I->selfDeclare( classFile );
			i2it++;
		}//while
		Image3DMap::iterator i3it;
		i3it = image3Ds().begin();
		while( i3it != image3Ds().end() ){
			Image3D * I = i3it->second;
			I->selfDeclare( classFile );
			i3it++;
		}//while
	}

	void
		Pipeline::generateExternalStageDeclarations( ofstream * classFile )
	{
		*classFile << std::endl;
		*classFile << "// externally defined stages" << std::endl;
		StageMap::iterator stageIt;
		stageIt = stages().begin();
		while( stageIt != stages().end() ){
			Stage * S = stageIt->second;
			S->selfDeclareExternal( classFile );
			stageIt++;
		}//while

	}


	void
		Pipeline::generateSynchronizationStageDeclarations( ofstream * classFile )
	{
		StageMap::iterator stageIt;
		stageIt = stages().begin();
		bool first = true;
		std::set<Stage *> barrierStages;
		while( stageIt != stages().end() ){
			Stage * S = stageIt->second;
			if( S->isSynchronizationStage() ){
				bool found = false;
				for( std::set<Stage *>::iterator sIt = barrierStages.begin();
					sIt != barrierStages.end(); ++sIt ){
						Stage * SS = *sIt;
						if( SS->stringValue( "barrier" ).compare( S->stringValue( "barrier" ) ) == 0 ){
							found = true;
							break;
						}
				}//for
				if( !found ){
					barrierStages.insert( S );

				}
			}
			stageIt++;
		}//while

		if( barrierStages.size() > 0 ){
			*classFile << std::endl;
			*classFile << "// synchronization stages" << std::endl;
		}
		std::set<Stage *>::iterator sIt = barrierStages.begin();
		for( ; sIt != barrierStages.end(); sIt++ ){
			Stage * S = *sIt;
			S->declareSynchronizationObject( classFile );
		}//for
	}



	void
		Pipeline::generateExternalStageInvocations( ofstream * classFile )
	{
		*classFile << std::endl;
		StageMap::iterator stageIt;
		stageIt = stages().begin();
		while( stageIt != stages().end() ){
			Stage * S = stageIt->second;
			S->selfInvokeExternal( classFile );
			stageIt++;
		}//while

	}


	void
		Pipeline::generateUniformBuffers( ofstream * classFile )
	{
		*classFile << std::endl;
		KernelMap::iterator kIt;
		kIt = kernels().begin();
		while( kIt != kernels().end() ){
			Kernel * k = kIt->second;
			k->selfGenerateUniformBuffers( classFile );
			kIt++;
		}//while
	}

	void
		Pipeline::generateKernelDeclarations( ofstream * classFile )
	{
		*classFile << std::endl;
		*classFile << "// Kernels" << std::endl;

		KernelMap::iterator kernelIt;
		kernelIt = kernels().begin();
		while( kernelIt != kernels().end() ){
			Kernel * K = kernelIt->second;
			K->selfDeclare( classFile );
			kernelIt++;
		}//while

	}


	void
		Pipeline::generateUniformDeclarations( ofstream * classFile )
	{
		*classFile << std::endl;
		*classFile << "// Uniforms" << std::endl;
		UniformMap::iterator unIt;
		unIt = uniforms().begin();
		while( unIt != uniforms().end() ){
			Uniform * U = unIt->second;
			U->selfDeclare( classFile );
			unIt++;
		}//while
	}


	void
		Pipeline::generateAccessorMutators( ofstream * classFile )
	{
		*classFile << std::endl;
		EquationMap::iterator eqIt;
		std::set<Uniform *> seen;
		eqIt = equations().begin();
		while( eqIt != equations().end() ){
			Equation * E = eqIt->second;
			UniformMap::iterator unIt = uniforms().find( E->name() );
			if( unIt != uniforms().end() ){
				Uniform * U = unIt->second;
				seen.insert( U );
				if( U->type().compare( "string" )){
					E->selfDeclareAccessorMutator( classFile );
				}
			}
			eqIt++;
		}//while
		for( UniformMap::iterator unIt = uniforms().begin();
			unIt != uniforms().end(); ++unIt ){
				Uniform * U = unIt->second;
				if( seen.find( U ) == seen.end() &&
					U->type().compare( "string" ) ){
					U->selfDeclareAccessorMutator( classFile );
				}
		}//for
	}


	void
		Pipeline::generateAttributeDeclarations( ofstream * classFile )
	{
		*classFile << std::endl;
		*classFile << "// Pipeline attributes" << std::endl;
		AttributeMap::iterator atIt;
		atIt = attributes().begin();
		while( atIt != attributes().end() ){
			Attribute * a = atIt->second;
			if( a->isInstantiable() ){
				a->selfDeclare( classFile );
			}
			atIt++;
		}//while
	}


	void
		Pipeline::generateTestDeclarations( ofstream * classFile )
	{
		*classFile << std::endl;
		*classFile << "// Unit tests" << std::endl;
		*classFile << "#if !defined( NDEBUG )" << std::endl;
		*classFile << "bool testBegin( QueueType queue );" << std::endl;
		*classFile << "bool testEnd( QueueType queue );" << std::endl;
		StageMap::iterator stageIt;
		stageIt = stages().begin();
		while( stageIt != stages().end() ){
			Stage * S = stageIt->second;
			S->selfDeclareTest( classFile );
			stageIt++;
		}//while
		*classFile << "#endif//NDEBUG" << std::endl;
	}


	void
		Pipeline::generateInitializeBuffersFunction( ofstream * classFile )
	{
		*classFile << std::endl;
		*classFile << "void initializeBuffers( void ){" << std::endl;
		BufferMap::iterator bufIt;
		bufIt = buffers().begin();
		while( bufIt != buffers().end() ){
			Buffer * B = bufIt->second;
			B->selfInitialize( classFile );
			bufIt++;
		}//while
		Image2DMap::iterator i2it;
		i2it = image2Ds().begin();
		while( i2it != image2Ds().end() ){
			Image2D * I = i2it->second;
			I->selfInitialize( classFile );
			i2it++;
		}//while
		Image3DMap::iterator i3it;
		i3it = image3Ds().begin();
		while( i3it != image3Ds().end() ){
			Image3D * I = i3it->second;
			I->selfInitialize( classFile );
			i3it++;
		}//while
		*classFile << "}" << std::endl;
	}


	void
		Pipeline::generateDirectComputeBufferDeclarations( ofstream * classFile )
	{
		if( targetIsDirectCompute() ){
			*classFile << std::endl;
			StageMap::iterator stIt;
			stIt = stages().begin();
			while( stIt != stages().end() ){
				Stage * s = stIt->second;
				s->selfDeclareDirectComputeBuffers( classFile );
				stIt++;
			}//while
		}
	}


	void
		Pipeline::generateInitializeStagesFunction( ofstream * classFile )
	{
		*classFile << std::endl;
		*classFile << "void initializeStages( void ){" << std::endl;
		StageMap::iterator stIt;
		stIt = stages().begin();
		while( stIt != stages().end() ){
			Stage * s = stIt->second;
			if( targetIsDirectCompute() ){
				s->selfInitializeUniformBuffers( classFile );
			}
			if( s->isSynchronizationStage() ){
				*classFile << "_initialize_sync_" << s->stringValue( "barrier" ) << "();" << std::endl;
			}
			stIt++;
		}//while
		*classFile << "}" << std::endl;
	}


	void
		Pipeline::generateInitializeKernelsFunction( ofstream * classFile )
	{
		*classFile << std::endl;
		*classFile << "void initializeKernels( void ){" << std::endl;

		KernelMap::iterator kernelIt;
		kernelIt = kernels().begin();
		while( kernelIt != kernels().end() ){
			Kernel * k = kernelIt->second;
			k->selfInitialize( classFile );
			kernelIt++;
		}//while

		*classFile << "}" << std::endl;
	}


	void
		Pipeline::generateInitializeAttributesFunction( ofstream * classFile )
	{
		*classFile << std::endl;
		*classFile << "void initializeAttributes( void ){" << std::endl;
		AttributeMap::iterator aIt;
		aIt = attributes().begin();
		while( aIt != attributes().end() ){
			Attribute * a = aIt->second;
			if( a->isInstantiable() ){
				ASTExpr v = a->value();
				*classFile << a->name() << " = " << stringify( v ) << ";" << std::endl;
			}
			aIt++;
		}//while
		*classFile << "}" << std::endl;
	}


	void
		Pipeline::generateEquationEvaluationFunctions( ofstream * classFile )
	{
		EquationMap::iterator eqIt;
		eqIt = equations().begin();
		while( eqIt != equations().end() ){
			Equation * e = eqIt->second;
			e->generateSelfEvaluation( classFile );
			eqIt++;
		}//while
	}


	void
		Pipeline::generateInitializeEquationsFunction( ofstream * classFile )
	{
		*classFile << std::endl;
		*classFile << "void initializeEquations( void ){" << std::endl;

		bool done = false;
		while( !done ){
			done = true;
			EquationMap::iterator eqIt;
			eqIt = equations().begin();
			while( eqIt != equations().end() ){
				Equation * e = eqIt->second;
				eqIt++;
				if( e->evaluated()){
					continue;
				}
				if( e->allPrecedentsEvaluated() ){
					e->generateSelfEvaluationAssignment( classFile );
					e->evaluatedIs( true );
					done = false;
				}
				else{
					if( e->undefined() && !e->evaluated() ){
						*classFile << "// defer " << e->name() << " due to undefined precedents" << std::endl;
						e->evaluatedIs( true );
					}
				}
			}//while
		}//while

		*classFile << "}" << std::endl;
	}


	void
		Pipeline::generateInitializeFunction( ofstream * classFile )
	{
		generateInitializeAttributesFunction( classFile );
		generateEquationEvaluationFunctions( classFile );
		generateInitializeEquationsFunction( classFile );
		generateInitializeBuffersFunction( classFile );
		generateInitializeKernelsFunction( classFile );
		generateDirectComputeBufferDeclarations( classFile );
		generateInitializeStagesFunction( classFile );
		*classFile << std::endl;
		*classFile << "void initializeScalars(){" << std::endl;
		*classFile << "initializeAttributes();" << std::endl;
		*classFile << "initializeEquations();" << std::endl;
		*classFile << "}" << std::endl;
		*classFile << std::endl;
		*classFile << "void initializeStructures(){" << std::endl;
		*classFile << "initializeBuffers();" << std::endl;
		*classFile << "initializeKernels();" << std::endl;
		*classFile << "initializeStages();" << std::endl;
		*classFile << "}" << std::endl;
		*classFile << std::endl;
		*classFile << "void initialize( void ){" << std::endl;
		*classFile << "initializeScalars();" << std::endl;
		*classFile << "initializeStructures();" << std::endl;
		*classFile << "}" << std::endl;
	}


	void
		Pipeline::generateTerminateFunction( ofstream * classFile )
	{
		*classFile << std::endl;
		*classFile << "void terminate( void ){" << std::endl;

		*classFile << "}" << std::endl;
	}


	bool
		Pipeline::hasBufferAttribute( std::string attribute )
	{
		BufferMap::iterator bufIt;
		bufIt = buffers().begin();
		while( bufIt != buffers().end() ){
			Buffer * B = bufIt->second;
			if( B->isTrue( attribute )){
				return true;
			}
			bufIt++;
		}//while
		return false;
	}


	bool 
		Pipeline::hasGlInteropBuffer()
	{
		return hasBufferAttribute( "glInterop" );
	}


	bool
		Pipeline::hasDxInteropBuffer()
	{
		return hasBufferAttribute( "dxInterop" );
	}

	bool Pipeline::targetIsOpenCL()
	{
		Maybe<ASTExpr>::MaybeType me = maybeExpr( "target" );
		if( !Maybe<ASTExpr>::toBool( me )){
			return true;//default, undefined
		}
		ASTExpr v = Maybe<ASTExpr>::toJust(  me );
		std::string target = stringify( v );
		if( 0 == target.compare( "OpenCL" )){
			return true;
		}
		return false;
	}


	bool Pipeline::targetIsDirectCompute()
	{
		Maybe<ASTExpr>::MaybeType me = maybeExpr( "target" );
		if( !Maybe<ASTExpr>::toBool( me )){
			return false;//default, undefined
		}
		ASTExpr v = Maybe<ASTExpr>::toJust(  me );
		std::string target = stringify( v );
		if( 0 == target.compare( "DirectCompute" )){
			return true;
		}
		return false;
	}


	void
		Pipeline::invokeAcquireInterop( ofstream * classFile )
	{
		*classFile << "acquireInterop();" << std::endl;
	}


	void
		Pipeline::invokeReleaseInterop( ofstream * classFile )
	{
		*classFile << "releaseInterop();" << std::endl;
	}


	void
		Pipeline::generateAcquireInterop( ofstream * classFile )
	{
		*classFile << std::endl;
		*classFile << "void acquireInterop(){" << std::endl;

		if( hasGlInteropBuffer() && targetIsOpenCL() ){
			*classFile << "#if defined( GL_PIPELINE )" << std::endl;
			*classFile << "for( unsigned int i = 0; i < devices.size(); ++i ){" << std::endl;
			*classFile << "std::vector<cl::Memory> glObjects;" << std::endl;
			BufferMap::iterator bufIt;
			bufIt = buffers().begin();
			while( bufIt != buffers().end() ){
				Buffer * B = bufIt->second;
				if( B->isTrue( "glInterop" )){
					*classFile << "glObjects.push_back( " << B->name() << " );" << std::endl;
				}
				bufIt++;
			}//while
			*classFile << "queues[ i ].enqueueAcquireGLObjects( &glObjects );" << std::endl;
			*classFile << "queues[ i ].finish();" << std::endl;
			*classFile << "}//for" << std::endl;
			*classFile << "#endif" << std::endl;
		}
		if( hasDxInteropBuffer() && targetIsOpenCL() ){
			*classFile << "#if defined( DX_PIPELINE )" << std::endl;
			*classFile << "for( unsigned int i = 0; i < devices.size(); ++i ){" << std::endl;
			*classFile << "std::vector<cl::Memory> dxObjects;" << std::endl;
			BufferMap::iterator bufIt;
			bufIt = buffers().begin();
			while( bufIt != buffers().end() ){
				Buffer * B = bufIt->second;
				if( B->isTrue( "dxInterop" )){
					*classFile << "dxObjects.push_back( " << B->name() << " );" << std::endl;
				}
				bufIt++;
			}//while
			*classFile << "queues[ i ].enqueueAcquireD3D10Objects( &dxObjects );" << std::endl;
			*classFile << "queues[ i ].finish();" << std::endl;
			*classFile << "}//for" << std::endl;
			*classFile << "#endif" << std::endl;
		}
		*classFile << "}" << std::endl;
	}


	void
		Pipeline::generateReleaseInterop( ofstream * classFile )
	{
		*classFile << std::endl;
		*classFile << "void releaseInterop(){" << std::endl;

		if( hasGlInteropBuffer() && targetIsOpenCL() ){
			*classFile << "#if defined( GL_PIPELINE )" << std::endl;
			*classFile << "for( unsigned int i = 0; i < devices.size(); ++i ){" << std::endl;
			*classFile << "std::vector<cl::Memory> glObjects;" << std::endl;
			BufferMap::iterator bufIt;
			bufIt = buffers().begin();
			while( bufIt != buffers().end() ){
				Buffer * B = bufIt->second;
				if( B->isTrue( "glInterop" )){
					*classFile << "glObjects.push_back( " << B->name() << " );" << std::endl;
				}
				bufIt++;
			}//while
			*classFile << "queues[ i ].enqueueReleaseGLObjects( &glObjects );" << std::endl;
			*classFile << "queues[ i ].finish();" << std::endl;
			*classFile << "}//for" << std::endl;
			*classFile << "#endif" << std::endl;
		}
		if( hasDxInteropBuffer() && targetIsOpenCL() ){
			*classFile << "#if defined( DX_PIPELINE )" << std::endl;
			*classFile << "for( unsigned int i = 0; i < devices.size(); ++i ){" << std::endl;
			*classFile << "std::vector<cl::Memory> dxObjects;" << std::endl;
			BufferMap::iterator bufIt;
			bufIt = buffers().begin();
			while( bufIt != buffers().end() ){
				Buffer * B = bufIt->second;
				if( B->isTrue( "dxInterop" )){
					*classFile << "dxObjects.push_back( " << B->name() << " );" << std::endl;
				}
				bufIt++;
			}//while
			*classFile << "queues[ i ].enqueueReleaseD3D10Objects( &dxObjects );" << std::endl;
			*classFile << "queues[ i ].finish();" << std::endl;
			*classFile << "}//for" << std::endl;
			*classFile << "#endif" << std::endl;
		}

		*classFile << "}" << std::endl;
	}


	void
		Pipeline::generateStageInvocations( ofstream * classFile )
	{
		*classFile << std::endl;
		*classFile << "// Stage invocations" << std::endl;

		StageMap::iterator stageIt;
		stageIt = stages().begin();
		while( stageIt != stages().end() ){
			Stage * S = stageIt->second;
			S->selfInvokeRun( classFile );
			S->selfInvokeTest( classFile );
			S->selfInvokeCapture( classFile );
			stageIt++;
		}//while
	}



	void
		Pipeline::generateDerivedGraphFunctions( ofstream * classFile )
	{
		{
			*classFile << std::endl;
			char * preamble = (char *)"void\ntraverseDerivedClassControlGraph(\n"
				"__SIGNATURE__,\n__PIPELINEMEMBER__ block\n)\n{\n";
			*classFile << preamble;
			*classFile << "(this->*block)( queue, e, success );" << std::endl;
			*classFile << "}\n" << std::endl;
		}
	}


	void
		Pipeline::generateGraphBasicBlocks( ofstream * classFile )
	{
		GraphMap::iterator gIt;
		gIt = graphs().begin();
		while( gIt != graphs().end() ){
			Graph * G = gIt->second;
			std::string startName = G->stringValue( "start" );
			GraphVertex * startVertex = G->vertices()[ startName ];
			G->clearState();
			G->generateBasicBlocks( classFile, startVertex, "_run" );
			*classFile << std::endl;
			*classFile << "#ifndef NDEBUG" << std::endl;
			G->clearState();
			G->generateBasicBlocks( classFile, startVertex, "_test" );
			G->clearState();
			G->generateBasicBlocks( classFile, startVertex, "_capture" );
			*classFile << "#endif//NDEBUG" << std::endl;
			gIt++;
		}//while
	}


	void
		Pipeline::generateClearSynchronization(
		ofstream * classFile
		)
	{
		*classFile << std::endl;
		*classFile << "void clearSynchronization(){" << std::endl;

		GraphMap::iterator gIt;
		gIt = graphs().begin();
		while( gIt != graphs().end() ){
			Graph * G = gIt->second;
			GraphVertexMap::iterator gvmIt;
			gvmIt = G->vertices().begin();
			while( gvmIt != G->vertices().end() ){
				GraphVertex * GV = gvmIt->second;
				Stage * S = GV->stage();
				if( S->isSynchronizationStage() ){
					*classFile << "clearSynchronizationObject( _sync_" << S->stringValue( "barrier" )
						<< ", _mutex_" << S->stringValue( "barrier" ) << " );" << std::endl;
				}
				gvmIt++;
			}//while
			gIt++;
		}//while

		*classFile << "}" << std::endl;

		*classFile << std::endl;
		*classFile << "void clearSynchronization( std::string graphName ){" << std::endl;

		gIt = graphs().begin();
		while( gIt != graphs().end() ){
			Graph * G = gIt->second;
			*classFile << "if( 0 == graphName.compare( \"" << G->name() << "\" )){" << std::endl;
			GraphVertexMap::iterator gvmIt;
			gvmIt = G->vertices().begin();
			while( gvmIt != G->vertices().end() ){
				GraphVertex * GV = gvmIt->second;
				Stage * S = GV->stage();
				if( S->isSynchronizationStage() ){
					*classFile << "clearSynchronizationObject( _sync_" << S->stringValue( "barrier" )
						<< ", _mutex_" << S->stringValue( "barrier" ) << " );" << std::endl;
				}
				gvmIt++;
			}//while
			*classFile << "}" << std::endl;
			gIt++;
		}//while

		*classFile << "}" << std::endl;
	}


	void
		Pipeline::generateStartSynchronization(
		ofstream * classFile,
		Graph * G
		)
	{
		GraphVertexMap::iterator gvmIt;
		gvmIt = G->vertices().begin();
		while( gvmIt != G->vertices().end() ){
			GraphVertex * GV = gvmIt->second;
			Stage * S = GV->stage();
			if( S->isSynchronizationStage() ){
				*classFile << "startSynchronization( _sync_" << S->stringValue( "barrier" )
					<< ", _mutex_" << S->stringValue( "barrier" ) << " );" << std::endl;
			}
			gvmIt++;
		}//while
	}


	void
		Pipeline::generateExecutionFunctions(
		std::string prefix,
		ofstream * classFile
		)
	{
		// with no arguments, invoke first alphebtic graph
		// else look up graph name and match to graph block, enqueue graph block, call acquirethread
		// use run/test/capture to enqueue correct block

		// ( block )

		*classFile << std::endl;
		*classFile << "void " << prefix << "ToCompletion( bool block = false ){" << std::endl;

		GraphMap::iterator grIt;
		grIt = graphs().begin();
		Graph * G = grIt->second;

		invokeAcquireInterop( classFile );

		if( targetIsOpenCL() ){
			*classFile << "for( unsigned int i = 0; i < queues.size(); ++i ){" << std::endl;
			*classFile << "QueueType queue = queues[ i ];" << std::endl;
		}

		generateStartSynchronization( classFile, G );

		if( 0 == prefix.compare( "test" ) ){
			*classFile << "bool allTestsPassed;" << std::endl;
			*classFile << "allTestsPassed = testBegin( queue );" << std::endl;
		}

		G->generateExecute( classFile, prefix );

		*classFile << "if( block ) while( isRunning() ){ SLEEP; }" << std::endl;
		if( 0 == prefix.compare( "test" ) ){
			*classFile << "allTestsPassed = testEnd( queue );" << std::endl;
		}

		if( targetIsOpenCL() ){
			*classFile << "}//for" << std::endl;
		}

		invokeReleaseInterop( classFile );

		*classFile << "}" << std::endl;

		// ( block, queue )

		*classFile << std::endl;
		*classFile << "void " << prefix << "ToCompletion( bool block, QueueType queue ){" << std::endl;

		invokeAcquireInterop( classFile );

		if( 0 == prefix.compare( "test" ) ){
			*classFile << "bool allTestsPassed;" << std::endl;
			*classFile << "allTestsPassed = testBegin( queue );" << std::endl;
		}

		grIt = graphs().begin();
		G = grIt->second;

		generateStartSynchronization( classFile, G );

		G->generateExecute( classFile, prefix );

		*classFile << "if( block ) while( isRunning() ){ SLEEP; }" << std::endl;
		if( 0 == prefix.compare( "test" ) ){
			*classFile << "allTestsPassed = testEnd( queue );" << std::endl;
		}

		invokeReleaseInterop( classFile );

		*classFile << "}" << std::endl;

		// ( block, graphName )

		*classFile << std::endl;
		*classFile << "void " << prefix << "ToCompletion( bool block, std::string graphName ){" << std::endl;

		invokeAcquireInterop( classFile );

		if( targetIsOpenCL() ){
			*classFile << "for( unsigned int i = 0; i < queues.size(); ++i ){" << std::endl;
			*classFile << "QueueType queue = queues[ i ];" << std::endl;
		}
		if( 0 == prefix.compare( "test" ) ){
			*classFile << "bool allTestsPassed;" << std::endl;
			*classFile << "allTestsPassed = testBegin( queue );" << std::endl;
		}

		while( grIt != graphs().end() ){
			Graph * G = grIt->second;
			*classFile << "if( 0 == graphName.compare( \"" << G->name() << "\" )){" << std::endl;

			generateStartSynchronization( classFile, G );

			G->generateExecute( classFile, prefix );

			*classFile << "}" <<std::endl;
			grIt++;
		}//while

		*classFile << "if( block ) while( isRunning() ){ SLEEP; }" << std::endl;
		if( 0 == prefix.compare( "test" ) ){
			*classFile << "allTestsPassed = testEnd( queue );" << std::endl;
		}

		if( targetIsOpenCL() ){
			*classFile << "}//for" << std::endl;
		}

		invokeReleaseInterop( classFile );

		*classFile << "}" << std::endl;

		// ( block, graphName, queue )

		*classFile << std::endl;
		*classFile << "void " << prefix << "ToCompletion( bool block, std::string graphName, QueueType queue ){" << std::endl;

		invokeAcquireInterop( classFile );

		if( 0 == prefix.compare( "test" ) ){
			*classFile << "bool allTestsPassed;" << std::endl;
			*classFile << "allTestsPassed = testBegin( queue );" << std::endl;
		}

		grIt = graphs().begin();
		while( grIt != graphs().end() ){
			Graph * G = grIt->second;
			*classFile << "if( 0 == graphName.compare( \"" << G->name() << "\" )){" << std::endl;

			generateStartSynchronization( classFile, G );

			G->generateExecute( classFile, prefix );

			*classFile << "}" <<std::endl;
			grIt++;
		}//while

		*classFile << "if( block ) while( isRunning() ){ SLEEP; }" << std::endl;
		if( 0 == prefix.compare( "test" ) ){
			*classFile << "allTestsPassed = testEnd( queue );" << std::endl;
		}

		invokeReleaseInterop( classFile );

		*classFile << "}" << std::endl;

	}


	void
		Pipeline::generateExecutionFunctions( ofstream * classFile )
	{
		generateExecutionFunctions( "run", classFile );
		*classFile << std::endl;
		*classFile << "#if !defined( NDEBUG )" << std::endl;
		generateExecutionFunctions( "test", classFile );
		generateExecutionFunctions( "capture", classFile );
		*classFile << std::endl;
		*classFile << "#endif" << std::endl;
	}


	void
		Pipeline::declareComputeShaderSourceStrings( ofstream * classFile )
	{
		KernelMap::iterator kIt;
		kIt = kernels().begin();
		while( kIt != kernels().end() ){
			Kernel * k = kIt->second;
			if( k->isDefined( "sourceFunction" )){
				Maybe<ASTExpr>::MaybeType me = k->maybeExpr( "sourceFunction" );
				ASTExpr v = Maybe<ASTExpr>::toJust( me );
				std::string sourceStringText = stringify( v );
				*classFile << "std::string " << sourceStringText << "();" << std::endl;
			}
			kIt++;
		}//while
	}

	void
		Pipeline::declareDirectComputeBuffers( Kernel * k, ofstream * kernelFile )
	{
		*kernelFile << std::endl;
		int idUAV = 0;
		int idSRV = 0;
		for( KernelParameterList::iterator kplIt = k->kernelParameters().begin();
			kplIt != k->kernelParameters().end(); kplIt++ ){
				KernelParameter kernelParameter = *kplIt;
				Buffer * b = dynamic_cast<Buffer *>( kernelParameter.node );
				b->selfDeclareDirectComputeBuffer( kernelFile, k, idUAV, idSRV );			
		}//for
#if 0
		BufferMap::iterator bufIt;
		bufIt = k->inputs().begin();
		int idUAV = 0;
		int idSRV = 0;
		while( bufIt != k->inputs().end() ){
			Buffer * b = bufIt->second;
			b->selfDeclareDirectComputeBuffer( kernelFile, k, idUAV, idSRV );			
			bufIt++;
		}//while
		bufIt = k->outputs().begin();
		while( bufIt != k->outputs().end() ){
			Buffer * b = bufIt->second;
			if( k->inputs().find( b->name() ) == k->inputs().end() ){
				b->selfDeclareDirectComputeBuffer( kernelFile, k, idUAV, idSRV );			
			}
			bufIt++;
		}//while
#endif
	}

	void
		Pipeline::declareConstantBuffers( Kernel * k, ofstream * kernelFile )
	{
		*kernelFile << std::endl;
		*kernelFile << "cbuffer UniformBuffer" << k->name() << " : register( c0 ){" << std::endl;
#if 0
		UniformMap::iterator unIt;
		unIt = k->uniforms().begin();
		int totalSize = 0;
		while( unIt != k->uniforms().end() ){
			Uniform * u = unIt->second;
			u->selfDeclare( kernelFile );
			totalSize += u->size();
			unIt++;
		}//while
#else
		int totalSize= 0;
		for( KernelParameterList::iterator kplIt = k->kernelParameters().begin();
			kplIt != k->kernelParameters().end(); ++kplIt ){
				KernelParameter kernelParameter = *kplIt;
				Node * node = kernelParameter.node;
				Uniform * u = dynamic_cast<Uniform *>( node );
				if( u != NULL ){
					u->selfDeclare( kernelFile );
					totalSize += u->size();
				}
		}//for
#endif

		int paddedSize = ( ( ( totalSize - 1 ) / 16 ) + 1 ) * 16;
		int padding = ( paddedSize - totalSize ) / sizeof( int );
		if( padding > 0 ){
			*kernelFile << "int __padding[ " << padding << " ];" << std::endl;
		}
		*kernelFile << "};" << std::endl;
		*kernelFile << std::endl;
	}


	void
		Pipeline::generateComputeShaderHeaders()
	{
		ofstream * shaderFile;
		KernelMap::iterator kIt;
		kIt = kernels().begin();
		while( kIt != kernels().end() ){
			Kernel * k = kIt->second;
			shaderFile = outputFile( k->name().c_str(), ".hlsl" );
			declareDirectComputeBuffers( k, shaderFile );
			declareConstantBuffers( k, shaderFile );
			k->selfGenerateHeader( shaderFile );
			shaderFile->close();
			kIt++;
		}//while
	}

	void
		Pipeline::generateKernelHeaders()
	{
		ofstream * kernelFile;
		kernelFile = outputFile( ".cl" );
		*kernelFile << std::endl;
		std::string dateTime = string( "date-time-tbd" );
		*kernelFile << "// This file was generated on " << dateTime << " by " 
			<< APPLICATION_NAME << " ver. " << APPLICATION_VERSION << std::endl;
		generateCLCodes( kernelFile );
		KernelMap::iterator kIt;
		kIt = kernels().begin();
		while( kIt != kernels().end() ){
			Kernel * k = kIt->second;
			k->selfGenerateHeader( kernelFile );
			kIt++;
		}//while
	}

	void
		Pipeline::checkName( std::string name )
	{
		std::string reservedWords[] = {
			"auto",
			"const",
			"double",
			"float",
			"int",
			"short",
			"struct",
			"unsigned",
			"break",
			"continue",
			"else",
			"for",
			"long",
			"signed",
			"switch",
			"void",
			"case",
			"default",
			"enum",
			"goto",
			"register",
			"sizeof",
			"typedef",
			"volatile",
			"char",
			"do",
			"extern",
			"if",
			"return",
			"static",
			"union",
			"while",
			"asm",
			"dynamic_cast",
			"namespace",
			"reinterpret_cast",
			"try",
			"bool",
			"explicit",
			"new",
			"static_cast",
			"typeid",
			"catch",
			"false",
			"operator",
			"template",
			"typename",
			"class",
			"friend",
			"private",
			"this",
			"using",
			"const_cast",  
			"inline",        
			"public",     
			"throw",             
			"virtual",
			"delete",      
			"mutable",       
			"protected",  
			"true",              
			"wchar_t",
			"and",      
			"bitand",   
			"compl",   
			"not_eq",   
			"or_eq",   
			"xor_eq",
			"and_eq",   
			"bitor",    
			"not",     
			"or",       
			"xor"
		};
		unsigned int numReservedWords = sizeof( reservedWords ) / sizeof( std::string );
		for( unsigned int i = 0; i < numReservedWords; ++i ){
			if( !name.compare( reservedWords[ i ] )){
				throw std::runtime_error(
					"invalid use of C++ reserved word '" + name + "' as identifier" );
			}
		}//for
	}

	void
		Pipeline::checkExpression( std::string lhsName, ASTExpr expression )
	{
		std::list< std::string > nameList;
		nameList = expr2NameList( expression );
		if( !( 1 == nameList.size() && 0 == nameList.front().compare( "undefined" ))){
			for( unsigned int i = 0; i < nameList.size(); ++i ){
				std::string name = nameList.front();
				nameList.pop_front();
				checkName( name );
				EquationMap::iterator eqIt = equations().find( name );
				Equation * equation = eqIt->second;
				bool sawEq = eqIt != equations().end() || isDefined( name );
				UniformMap::iterator unIt = uniforms().find( name );
				bool sawUn = unIt != uniforms().end();
				if( !sawEq && !sawUn && ( equation->type().compare( "string" ))){
					throw std::runtime_error( 
						"undefined term " + name + " in expression for "
						+ lhsName );
				}
			}//for
		}
	}

	void
		Pipeline::checkEquations()
	{
		EquationMap::iterator eqIt = equations().begin();
		while( eqIt != equations().end() ){
			Equation * E = eqIt->second;
			E->check();
			eqIt++;
		}//while
	}

	void 
		Pipeline::checkUniforms()
	{
		UniformMap::iterator unIt = uniforms().begin();
		while( unIt != uniforms().end() ){
			Uniform * U = unIt->second;
			U->check();
			unIt++;
		}//while
	}

	void 
		Pipeline::checkBuffers()
	{
		BufferMap::iterator bIt = buffers().begin();
		while( bIt != buffers().end() ){
			Buffer * B = bIt->second;
			B->check();
			bIt++;
		}//while
	}

	void 
		Pipeline::checkKernels()
	{
		KernelMap::iterator kIt = kernels().begin();
		while( kIt != kernels().end() ){
			Kernel * K = kIt->second;
			K->check();
			kIt++;
		}//while
	}

	void 
		Pipeline::checkStages()
	{
		StageMap::iterator sIt = stages().begin();
		while( sIt != stages().end() ){
			Stage * S = sIt->second;
			S->check();
			sIt++;
		}//while
	}

	void 
		Pipeline::checkGraphs()
	{
		GraphMap::iterator gIt = graphs().begin();
		while( gIt != graphs().end() ){
			Graph * G = gIt->second;
			G->check();
			gIt++;
		}//while
	}

	void
		Pipeline::checkStaticSemantics()
	{
		checkEquations();
		checkUniforms();
		checkBuffers();
		checkKernels();
		checkStages();
		checkGraphs();
	}


	void
		Pipeline::generatePipelineClass()
	{
		ofstream * classFile = outputFile( ".hpp" );
		*classFile << std::endl;
		*classFile << "#ifndef __" << name() << "_HPP__" << std::endl;
		*classFile << "#define __" << name() << "_HPP__" << std::endl;
		*classFile << std::endl;
		std::string dateTime = string( "date-time-tbd" );
		*classFile << "// This file was generated on " << dateTime << " by " 
			<< APPLICATION_NAME << " ver. " << APPLICATION_VERSION << std::endl;

		*classFile << "// Please do not edit this file.  Instead edit " << name()
			<< ".pl and rerun the " << APPLICATION_NAME << " tool." << std::endl;

		*classFile << std::endl;
		*classFile << "#ifdef __PIPELINE__" << std::endl;
		*classFile << "#undef __PIPELINE__" << std::endl;
		*classFile << "#undef __PIPELINEMEMBER__" << std::endl;
		*classFile << "#undef __RTCLASS__" << std::endl;
		*classFile << "#endif" << std::endl;

		*classFile << "#define __PIPELINE__ " << name() << std::endl;
		*classFile << "#define __PIPELINEMEMBER__ __p_" << name() << std::endl;
		*classFile << "#define __RTCLASS__ __r_" << name() << std::endl;

		if( targetIsOpenCL() ){
			*classFile << "#ifndef TARGET_OPENCL" << std::endl;
			*classFile << "#define TARGET_OPENCL" << std::endl;
			*classFile << "#endif" << std::endl;
		}
		else if( targetIsDirectCompute() ){
			*classFile << "#ifndef TARGET_DIRECTCOMPUTE" << std::endl;
			*classFile << "#define TARGET_DIRECTCOMPUTE" << std::endl;
			*classFile << "#endif" << std::endl;
		}

		bool glInterop;
		bool dxInterop;
		getInteropUsage( glInterop, dxInterop );
		if( dxInterop ){
			*classFile << "#ifdef DX_PIPELINE" << std::endl;
			*classFile << "#define USE_DX_INTEROP" << std::endl;
			*classFile << "#endif" << std::endl;
		}

		// RuntimePipeline class

		*classFile << boilerplate();

		// Top level user code

		generateCPPCodes( classFile, true );

		// kernel string declarations
		if( targetIsDirectCompute() ){
			declareComputeShaderSourceStrings( classFile );
		}

		// Class definition

		*classFile << std::endl;
		*classFile << "// " << name() << " class" << std::endl;
		*classFile << std::endl;
		*classFile << "class " << name() << " : public __RTCLASS__{" << std::endl;
		*classFile << "public:" << std::endl;

		if( targetIsOpenCL() ){
			*classFile << "std::string clSourceFileName( void ){";
			*classFile << "return \"" << name() << ".cl\"; }" << std::endl;
		}

		*classFile << "std::string pipelineName(){ return \"" << name() << "\"; }" << std::endl;

		generateCPPCodes( classFile, false );

		if( targetIsOpenCL() ){
			generateInteropBools( classFile );
		}

		generateAcquireInterop( classFile );
		generateReleaseInterop( classFile );
		generateBufferDeclarations( classFile );
		generateExternalStageDeclarations( classFile );
		generateExternalStageInvocations( classFile );
		generateSynchronizationStageDeclarations( classFile );
		generateUniformBuffers( classFile );
		generateKernelDeclarations( classFile );
		generateUniformDeclarations( classFile );
		generateAccessorMutators( classFile );
		generateAttributeDeclarations( classFile );
		generateTestDeclarations( classFile );
		generateInitializeFunction( classFile );
		generateTerminateFunction( classFile );
		generateStageInvocations( classFile );
		generateDerivedGraphFunctions( classFile );
		generateGraphBasicBlocks( classFile );
		generateClearSynchronization( classFile );
		generateExecutionFunctions( classFile );
		*classFile << "};" << std::endl;
		*classFile << "#endif//__" << name() << "_HPP__" << std::endl;
		classFile->close();
	}

	void
		Pipeline::generatePipelineMain()
	{
		ofstream * mainFile = outputFile( ".main.cpp" );
		if( mainFile != NULL ){
			*mainFile << std::endl;
			*mainFile << "#include \"" << name() << ".hpp\"" << std::endl;
			*mainFile << std::endl;
			*mainFile << "int main( int argc, char * argv[] ){" << std::endl;
			*mainFile << name() << " * " << name() << "_ = new " << name() << "();" << std::endl;
			*mainFile << name() << "_->initializeOpenCL();" << std::endl;
			*mainFile << name() << "_->initialize();" << std::endl;
			for( GraphMap::iterator gIt = graphs().begin(); gIt != graphs().end();  ++gIt ){
				Graph * graph = gIt->second;
				*mainFile << name() << "_->runToCompletion( true, \""
					<< graph->name() << "\" );" << std::endl;
			}//for
			*mainFile << "}" << std::endl;
			mainFile->close();
		}
	}


	void
		Pipeline::generatePipelineExternals()
	{
		ofstream * extFile = outputFile( ".ext.cpp" );
		*extFile << std::endl;
		*extFile << "// This file was generated on " << "date-time-tbd" << " by " 
			<< APPLICATION_NAME << " ver. " << APPLICATION_VERSION << std::endl;
		*extFile << "// Please edit this file to include invocations of external stages." << std::endl;
		*extFile << "#include \"" << name() << ".hpp\"" << std::endl;

		*extFile << std::endl;
		*extFile << "#ifdef TARGET_OPENCL" << std::endl;
		*extFile << "std::string __RTCLASS__::predefinedStrings( void )" << std::endl;
		*extFile << "{" << std::endl;
		*extFile << "#ifndef NDEBUG" << std::endl;
		*extFile << "return std::string( \"\" );" << std::endl;
		*extFile << "#else" << std::endl;
		*extFile << "return std::string( \"-DNDEBUG\" );" << std::endl;
		*extFile << "#endif" << std::endl;
		*extFile << "}" << std::endl;
		*extFile << "#endif" << std::endl;

		StageMap::iterator stageIt;
		stageIt = stages().begin();
		while( stageIt != stages().end() ){
			Stage * S = stageIt->second;
			if( S->isDeclaredExternally() ){
				S->selfImplementExternal( extFile );
			}
			stageIt++;
		}//while
		if( extFile != NULL ){
			extFile->close();
		}
	}

	void
		Pipeline::generatePipelineTest()
	{
		ofstream * testFile = outputFile( ".test.cpp" );
		*testFile << "// This file was generated on " << "date-time-tbd" << " by " 
			<< APPLICATION_NAME << " ver. " << APPLICATION_VERSION << std::endl;
		*testFile << "// Please edit this file to include unit tests for each stage." << std::endl;

		*testFile << std::endl;
		*testFile << "#include \"" << name() << ".hpp\"" << std::endl;

		*testFile << "#if !defined( NDEBUG )" << std::endl;

		*testFile << std::endl;
		*testFile << "bool" << std::endl;
		*testFile << name() << "::testBegin( QueueType queue )" << std::endl;
		*testFile << "{" << std::endl;
		*testFile << "// Please edit this to test preconditions for the pipeline" << std::endl;
		*testFile << "return true;" << std::endl;
		*testFile << "}" << std::endl;

		StageMap::iterator stageIt;
		stageIt = stages().begin();
		while( stageIt != stages().end() ){
			Stage * S = stageIt->second;
			S->selfImplementTest( testFile );
			stageIt++;
		}//while

		*testFile << std::endl;
		*testFile << "bool" << std::endl;
		*testFile << name() << "::testEnd( QueueType queue )" << std::endl;
		*testFile << "{" << std::endl;
		*testFile << "// Please edit this to test postconditions for the pipeline" << std::endl;
		*testFile << "return true;" << std::endl;
		*testFile << "}" << std::endl;

		*testFile << "#endif//NDEBUG" << std::endl;

		testFile->close();
	}


	void
		Pipeline::generatePipelineVisualization( void )
	{
		ofstream * dotFile = outputFile( ".dot" );

		*dotFile << "// Pipeline " << name() << std::endl;
		GraphMap::iterator gmIt;
		gmIt = graphs().begin();
		while( gmIt != graphs().end() ){
			Graph * G = gmIt->second;
			G->generateVisualization( dotFile );
			gmIt++;
		}//while
		dotFile->close();
	}

#if 0
	void
		Pipeline::findEquationTypes()
	{
		for( EquationMap::iterator eqIt = equations().begin(); eqIt != equations().end(); ++eqIt ){
			Equation * equation = eqIt->second;
			UniformMap::iterator unIt = uniforms().find( equation->name() );
			if( unIt != uniforms().end() ){
				Uniform * uniform = unIt->second;
				equation->typeIs( uniform->type() );
			}
		}//for
	}
#endif

	void
		Pipeline::generateRuntimePipeline()
	{
		checkStaticSemantics();
#if 0
		findEquationTypes();
#endif
		generatePipelineClass();
		if( targetIsOpenCL() ){
			generateKernelHeaders();
		}
		else if( targetIsDirectCompute() ){
			generateComputeShaderHeaders();
		}
		generatePipelineMain();
		generatePipelineExternals();
		generatePipelineTest();
		generatePipelineVisualization();
	}



}//namespace amd

